﻿////////////////////////////////////////////////////////////////
///
/// File: Translator.cs
/// Author: Cristian Dittamo
/// Last update: 25 March 2010
/// Description: This file contains the definition of the 
///              main classes needed for translating MSIL into nVidia PTX code.
/// To do: 
/// ------------------------------------------------------------
/// Copyright (c) 2009-2010 Cristian Dittamo (dittamo@di.unipi.it)
/// 
/// The use and distribution terms for this software are 
/// contained in the file named license.txt, which can be found 
/// in the root of this distribution.
/// By using this software in any fashion, you are agreeing to 
/// be bound by the terms of this license.
///
/// You must not remove this notice, or any other, from this
/// software.
/// ------------------------------------------------------------
////////////////////////////////////////////////////////////////

using System;
using System.Collections;
using System.Collections.Generic;
using System.Text;
using System.Runtime.InteropServices;
using System.ComponentModel;
using System.Reflection;

using MappedView;
using CLIFile;

using StreamDefs;

namespace CompileLib
{
    namespace Translation
    {
        using System.IO;
        using System.Reflection.Emit;
        using InteropServices;

        public enum MemoryAccessType
        {
            RO = 1,
            RW
        }

        public class MetaKernel
        {
            protected ILCursor _cursor;
            protected Object instance;
            protected int _cores;
            protected MethodInfo _meth;
            /// <summary>
            /// Methods invoked from the current kernel
            /// </summary>
            protected List<MethodInfo> _invokedMethods;
            /// <summary>
            /// Memory local to threads
            /// </summary>
            protected IList<LocalVariableInfo> _localMemory;
            protected CLIType[] _locvars;
            protected CLIType[] _args;
            protected CLIFileReader _clifile;
            protected SystemType sType;

            protected Dictionary<System.Reflection.Emit.OpCode, int> _registers;

            /// <summary>
            /// Shared memory between instances (threads) of the current kernel
            /// </summary>
            protected Dictionary<FieldInfo, MemoryAccessType> _sharedMemory;
            /// <summary>
            /// Shared memory between instances (threads) of the current kernel
            /// </summary>
            protected Dictionary<FieldInfo, MemoryAccessType> _globalMemory;
            /// <summary>
            /// Streams for current kernels
            /// </summary>
            protected ParameterInfo[] _inputs;

            public MetaKernel(Object worker, MethodInfo method, SystemType st)
            {
                instance = worker;

                sType = st;

                _meth = method;
                System.Reflection.MethodBody mb = _meth.GetMethodBody();
                /// Evaluate memory resources: local memory
                _localMemory = mb.LocalVariables;

                MethodTableCursor cur = CLIFileReader.FindMethod(method);
                CLIFile.MethodBody cb = cur.MethodBody;
                _cursor = cb.ILInstructions;

                //####################################################################################################

                LocalsVarSig locals = cb.LocalsSignature;
                if (locals != null)
                {
                    IEnumerator<CLIType> iterLocal = locals.GetVariables().GetEnumerator();
                    _locvars = new CLIType[locals.Count];
                    int l = 0;
                    while ((iterLocal.MoveNext()) && (l < _locvars.Length))
                        _locvars[l++] = iterLocal.Current;
                }

                MethodSig ms = cur.GetMethodSignature();
                _args = new CLIType[ms.Count + 1];
                IEnumerator<CLIType> iter = ms.GetParameters(cur.FindCLIType()).GetEnumerator();
                int m = 0;
                _args[m++] = cur.FindCLIType();
                while ((m < _args.Length) && (iter.MoveNext()))
                    _args[m++] = iter.Current;

                //###################################################################################################
                _inputs = _meth.GetParameters();
                //_outputs = null;

                _registers = new Dictionary<System.Reflection.Emit.OpCode, int>();

                _cores = 0;

                _sharedMemory = new Dictionary<FieldInfo, MemoryAccessType>();
                _globalMemory = new Dictionary<FieldInfo, MemoryAccessType>();
                _invokedMethods = new List<MethodInfo>();

                EvaluateResource();
            }
            
            public MetaKernel(Object worker, MethodInfo method, int cores, SystemType st)
            {
                instance = worker;

                sType = st;

                _meth = method;
                System.Reflection.MethodBody mb = _meth.GetMethodBody();
                /// Evaluate memory resources: local memory
                _localMemory = mb.LocalVariables;

                MethodTableCursor cur = CLIFileReader.FindMethod(method);
                CLIFile.MethodBody cb = cur.MethodBody;
                _cursor = cb.ILInstructions;

                //####################################################################################################

                 LocalsVarSig locals = cb.LocalsSignature;
                 if (locals != null)
                 {
                     IEnumerator<CLIType> iterLocal = locals.GetVariables().GetEnumerator();
                     _locvars = new CLIType[locals.Count];
                     int l = 0;
                     while ((iterLocal.MoveNext()) && (l < _locvars.Length))
                         _locvars[l++] = iterLocal.Current;
                 }
                
                MethodSig ms = cur.GetMethodSignature();
                _args = new CLIType[ms.Count + 1];
                IEnumerator<CLIType> iter = ms.GetParameters(cur.FindCLIType()).GetEnumerator();
                int m = 0;
                _args[m++] = cur.FindCLIType();
                while ((m < _args.Length) && (iter.MoveNext()))
                    _args[m++] = iter.Current;

                //###################################################################################################
                _inputs = _meth.GetParameters();
               
                _registers = new Dictionary<System.Reflection.Emit.OpCode, int>();

                _cores = cores;

                _sharedMemory = new Dictionary<FieldInfo, MemoryAccessType>();
                _globalMemory = new Dictionary<FieldInfo, MemoryAccessType>();
                _invokedMethods = new List<MethodInfo>();

                EvaluateResource();
            }

            public ILCursor Cursor
            {
                get { return _cursor; }
            }

            public int NumOfArguments()
            {
                return _inputs.Length;
            }

            public Type DeclaringType
            {
                get { return _meth.DeclaringType; }
            }

            public void EvaluateResource()
            {
                /// Evaluate kernel method body
                _cursor.TrackTargets();
                while (_cursor.Next())
                {
                    ILInstruction il = _cursor.Instr;
                    object parameter = il.ResolveParameter(_cursor.FileReader);
                    if (parameter != null)
                    {
                        Console.WriteLine("------MIL-- IL_{0} {1} {2}", _cursor.Position.ToString(), il.op.Name, parameter.ToString());
                        /// Evaluate memory resources: instance and static fields
                        EvalFieldOpCode(il, parameter);
                    }
                    else
                        Console.WriteLine("------MIL-- IL_{0} {1}", _cursor.Position.ToString(), il.op.Name);
                }

                //// register
                //foreach(LocalVariableInfo lvi in _localMemory)
                //{
                //    Console.WriteLine("LocalMem: {0}", lvi.LocalType.ToString());
                //    OperandType opt = TypesMapping.Convert(lvi.LocalType, sType);
                //   _ptxCode.AddRegister(new LocalMem(opt, String.Format("r{0}",lvi.LocalIndex),1));
                //}

                //// shared mem
                //Dictionary<FieldInfo, MemoryAccessType>.Enumerator shdIter = _sharedMemory.GetEnumerator();
                //while(shdIter.MoveNext())
                //{
                //    Console.WriteLine("SharedMem: {0}", shdIter.Current.Key.FieldType.ToString());
                //    OperandType opt = TypesMapping.Convert(shdIter.Current.Key.FieldType, sType);
                //    _ptxCode.AddSharedMem(new SharedMem(opt,0,String.Format("__cuda__sh{0}",shdIter.Current.Key.Name),1024));
                //}
            }

            /// <summary>
            /// Eval the current opcode
            /// TODO: Consider Ldflda/Stflda
            /// </summary>
            /// <param name="il">Current IL instruction</param>
            private void EvalFieldOpCode(ILInstruction il, object parameter)
            {
                switch ((OPCODES)(il.op.Value < 0 ? 256 + (il.op.Value & 0xFF) : il.op.Value))
                {
                    case OPCODES.Ldfld: //Shared memory
                        {
                            FieldInfo fi = ((FieldDesc)parameter).GetReflectionField();
                            if (!_sharedMemory.ContainsKey(fi))
                                _sharedMemory.Add(fi, MemoryAccessType.RO);
                        }
                        break;
                    case OPCODES.Stfld:  //Shared memory
                        {
                            FieldInfo fi = ((FieldDesc)parameter).GetReflectionField();
                            if (!_sharedMemory.ContainsKey(fi))
                                _sharedMemory.Add(fi, MemoryAccessType.RW);
                            else
                                _sharedMemory[fi] = MemoryAccessType.RW;
                        }
                        break;
                    case OPCODES.Ldsfld: //Global memory
                        {
                            FieldInfo fi = ((FieldDesc)parameter).GetReflectionField();
                            if (!_globalMemory.ContainsKey(fi))
                                _globalMemory.Add(fi, MemoryAccessType.RO);
                        }
                        break;
                    case OPCODES.Stsfld: //Global memory
                        {
                            FieldInfo fi = ((FieldDesc)parameter).GetReflectionField();
                            if (!_globalMemory.ContainsKey(fi))
                                _globalMemory.Add(fi, MemoryAccessType.RW);
                            else
                                _globalMemory[fi] = MemoryAccessType.RW;
                        }
                        break;
                }
            }

            private int evalNumPops()
            {
                int result = 0;
                switch (_cursor.Instr.op.StackBehaviourPop)
                {
                    case StackBehaviour.Pop0:
                        result = 0;
                        break;
                    case StackBehaviour.Pop1:
                    case StackBehaviour.Popi:
                    case StackBehaviour.Popref:
                        result = 1;
                        break;
                    case StackBehaviour.Pop1_pop1:
                    case StackBehaviour.Popi_pop1:
                    case StackBehaviour.Popi_popi8:
                    case StackBehaviour.Popi_popr4:
                    case StackBehaviour.Popi_popr8:
                    case StackBehaviour.Popref_pop1:
                    case StackBehaviour.Popref_popi:
                        result = 2;
                        break;
                    case StackBehaviour.Popi_popi_popi:
                    case StackBehaviour.Popref_popi_pop1:
                    case StackBehaviour.Popref_popi_popi:
                    case StackBehaviour.Popref_popi_popi8:
                    case StackBehaviour.Popref_popi_popr4:
                    case StackBehaviour.Popref_popi_popr8:
                    case StackBehaviour.Popref_popi_popref:
                        result = 3;
                        break;
                    case StackBehaviour.Varpop:
                        {
                            //int args = 0;
                            //if (_cursor.IsCall)
                            //{
                            //    object parameter = cil.ResolveParameter(_clifile);
                            //    MethodDesc m = (MethodDesc)parameter;
                            //    args = m.Signature.Count - 1 + (m.Signature.HasThis ? 1 : 0);
                            //}
                            //result = 1 + args;
                            result = 1;
                        }
                        break;
                }
                return result;
            }

            private int evalNumPushes()
            {
                int result = 0;
                switch (_cursor.Instr.op.StackBehaviourPush)
                {
                    case StackBehaviour.Push0:
                        result = 0;
                        break;
                    case StackBehaviour.Push1:
                    case StackBehaviour.Pushi:
                    case StackBehaviour.Pushi8:
                    case StackBehaviour.Pushr4:
                    case StackBehaviour.Pushr8:
                    case StackBehaviour.Pushref:
                        result = 1;
                        break;
                    case StackBehaviour.Varpush:
                        {
                            ILInstruction il = ILInstruction.Normalize(_cursor.Instr);
                            object res = il.ResolveParameter(_clifile);

                            if (res == null)
                                result = 0;
                            else
                            {
                                if (_cursor.Instr.op == System.Reflection.Emit.OpCodes.Call || _cursor.Instr.op == System.Reflection.Emit.OpCodes.Calli || _cursor.Instr.op == System.Reflection.Emit.OpCodes.Callvirt)
                                {
                                    MethodDesc m = (MethodDesc)res;
                                    CLIType ctp = m.Signature.ReturnType;
                                    Type tp = null;

                                    if( ctp.Type is VariableType )
                                        tp = ctp.GetReflectionType(new Type[] { res.GetType() });
                                    else
                                        tp = ctp.GetReflectionType();
                                    if (tp.Equals(typeof(void)))
                                        result = 0;
                                    else
                                        result = 1;
                                }
                                else
                                    result = 1;
                            }
                        }
                        break;
                    case StackBehaviour.Push1_push1:
                        result = 2;
                        break;
                }
                return result;
            }

            /// <summary>
            /// first element = number of pushes
            /// second element = number of pops
            /// </summary>
            /// <param name="op"></param>
            /// <returns></returns>
            protected int[] GetBehaviour(OPCODES op)
            {
                int[] access = new int[2];

                access[0] = evalNumPushes();

                if (_cursor.IsCall)
                    access[1] = _cursor.CallLoadArguments().Length;
                else
                {
                    if (op.CompareTo(OPCODES.Ret) == 0)
                    {
                        if (this._meth.ReturnType.Equals(typeof(void)))
                            access[1] = 0;
                        else
                            access[1] = evalNumPops();
                    }
                    else
                        access[1] = evalNumPops();
                }

                return access;
            }

            protected SystemType sysType
            {
                get { return sysType; }
            }
        }

        /// <summary>
        /// Manage _genCode for Kernels generation
        /// </summary>
        public class MetaCompute
        {
            protected Dictionary<string, MetaKernel> _kernels;
            protected FieldInfo[] _fields;
            protected CLIFileReader _clifile;
            protected int devices;
            
            protected Object instance;
            protected SystemType sType;

            public MetaCompute(Object worker, CLIFileReader clirdr, SystemType st)
            {
                instance = worker;
                _kernels = null;
                _clifile = clirdr;
                _kernels = new Dictionary<string, MetaKernel>();
                sType = st;
            }

            public virtual void AddKernel(MethodInfo kernel)
            {
                if (kernel == null)
                    throw new ArgumentOutOfRangeException();

                _kernels.Add(kernel.Name, new MetaKernel(instance, kernel, sType));
            }

            public MetaKernel GetKernel(string methodName)
            {
                return (MetaKernel)_kernels[methodName];
            }

            public SystemType SysType
            {
                get { return sType; }
                set { sType = value; }
            }
        }

        public enum SystemType
        {
            x64 = 0,
            x86 = 1
        }

        public static class MetaCompiler
        {
            /// <summary>
            /// Determine whether there is a GPU card installed in the current system or not.
            /// If yes, it should evaluate which GPU to use between ATI and NVIDIA.
            /// TODO: actually it statically selects NVIDIA 
            /// </summary>
            /// <param name="instanceType"></param>
            /// <param name="clifile"></param>
            /// <param name="compute"></param>
            /// <returns></returns>
            
            static public SystemType GetSystemType()
            {
                if (IntPtr.Size == 8)
                    return SystemType.x64;
                else
                    return SystemType.x86;
            }
        }

        namespace NVIDIA
        {
            using InteropServices.NVIDIA;

            #region PTX code modelling

            public class Predicate
            {
                string name;
                bool condition;
                string pred;

                public Predicate(bool cond, int id)
                {
                    StringBuilder tmp = new StringBuilder();
                    condition = cond;
                    name = String.Format("%p{0}", id);
                    if (condition)
                        tmp.AppendFormat("@{0}", name);
                    else
                        tmp.AppendFormat("@!{0}", name);
                    pred = tmp.ToString();
                }

                public string Name
                {
                    get { return name; }
                }

                public string CondToString()
                {
                    return pred;
                }

                public string DeclerToString()
                {
                    return ".pred " + name;
                }
            }

            public class OperandType : IComparable
            {
                // Type name
                String _name;

                // Type identifier : used for registers naming purpose
                String _id; 

                internal byte _s1;

                internal OperandType(string name, string id, byte s1)
                {
                    this._name = name;
                    this._s1 = s1;
                    this._id = id;
                }

                internal bool isNull()
                {
                    return (_name == null);
                }

                public override bool Equals(object obj)
                {
                    return ((obj is OperandType) && this.Equals((OperandType)obj));
                }

                public bool Equals(OperandType obj)
                {
                    return (obj._s1 == this._s1);
                }

                // Summary:
                //     Indicates whether two OpCode structures are not equal.
                //
                // Parameters:
                //   a:
                //     The OpCode to compare to b.
                //
                //   b:
                //     The OpCode to compare to a.
                //
                // Returns:
                //     true if a is not equal to b; otherwise, false.
                public static bool operator !=(OperandType a, OperandType b)
                {
                    return !(a.Equals(b));
                }

                //
                // Summary:
                //     Indicates whether two OpCode structures are equal.
                //
                // Parameters:
                //   a:
                //     The OpCode to compare to b.
                //
                //   b:
                //     The OpCode to compare to a.
                //
                // Returns:
                //     true if a is equal to b; otherwise, false.
                public static bool operator ==(OperandType a, OperandType b)
                {
                    return a.Equals(b);
                }

                public static bool operator >(OperandType a, OperandType b)
                {
                    return a._s1 > b._s1;
                }

                public static bool operator <(OperandType a, OperandType b)
                {
                    return a._s1 < b._s1;
                }

                public int CompareTo(object value)
                {
                    return this.Equals(value) ? 0 : 1;
                }

                //
                // Summary:
                //     The name of the NVIDIA PTX instruction.
                //
                // Returns:
                //     Read-only. The name of the NVIDIA PTX instruction.
                public string Name { get { return _name; } }

                public string Id4RegisterNaming { get { return _id; } }

                //
                // Summary:
                //     The operand type of an NVIDIA PTX instruction.
                //
                // Returns:
                //     Read-only. The operand type of an NVIDIA PTX instruction.
                //public OperandType OperandType { get { return _opt;} }
                public override int GetHashCode()
                {
                    return this._name.GetHashCode();
                }
            }

            public class OperandTypes
            {
                #region static fields
                public static readonly OperandType s8;
                public static readonly OperandType s16;
                public static readonly OperandType s32;
                public static readonly OperandType s64;
                public static readonly OperandType u8;
                public static readonly OperandType u16;
                public static readonly OperandType u32;
                public static readonly OperandType u64;
                public static readonly OperandType f16;
                public static readonly OperandType f32;
                public static readonly OperandType f64;
                public static readonly OperandType b8;
                public static readonly OperandType b16;
                public static readonly OperandType b32;
                public static readonly OperandType b64;
                public static readonly OperandType pred;
                public static readonly OperandType regname;

                public static readonly OperandType param;
                public static readonly OperandType target;
                public static readonly OperandType method;
                #endregion

                static OperandTypes()
                {
                    b8 = new OperandType("b8", "b", 0);
                    s8 = new OperandType("s8", "s", 1);
                    u8 = new OperandType("u8", "u", 2);
                    b16 = new OperandType("b16", "db", 3);
                    u16 = new OperandType("u16", "du", 4);
                    s16 = new OperandType("s16", "ds", 5);
                    f16 = new OperandType("f16", "df", 6);
                    b32 = new OperandType("b32", "tb", 7);
                    u32 = new OperandType("u32", "tu", 8);
                    s32 = new OperandType("s32", "ts", 9);
                    f32 = new OperandType("f32", "tf", 10);
                    b64 = new OperandType("b64", "qb", 11);
                    u64 = new OperandType("u64", "qu", 12);
                    s64 = new OperandType("s64", "qs", 13);
                    f64 = new OperandType("f64", "qf", 14);
                    pred = new OperandType("pred", "", 15);
                    param = new OperandType("param", "", 16);
                    target = new OperandType("target", "", 17);
                    method = new OperandType("method", "", 18);
                }
            }

            public class CompOp
            {
                private string op;
                internal CompOp(string o)
                {
                    op = o;
                }

                public override string ToString()
                {
                    return op;
                }
            }

            public class CompOps
            {
                //eq, ne, lt, le, gt, ge, lo, ls, hi, hs, equ, neu, ltu, leu, gtu, geu, num, nan 
                #region static fields
                public static readonly CompOp eq;
                public static readonly CompOp ne;
                public static readonly CompOp lt;
                public static readonly CompOp le;
                public static readonly CompOp gt;
                public static readonly CompOp ge;
                public static readonly CompOp lo;
                public static readonly CompOp ls;
                public static readonly CompOp hi;
                public static readonly CompOp hs;
                public static readonly CompOp equ;
                public static readonly CompOp neu;
                public static readonly CompOp ltu;
                public static readonly CompOp leu;
                public static readonly CompOp gtu;
                public static readonly CompOp geu;
                public static readonly CompOp num;
                public static readonly CompOp nan;
                #endregion

                static CompOps()
                {
                    eq = new CompOp("eq");
                    ne = new CompOp("ne");
                    lt = new CompOp("lt");
                    le = new CompOp("le");
                    gt = new CompOp("gt");
                    ge = new CompOp("ge");
                    lo = new CompOp("lo");
                    ls = new CompOp("ls");
                    hi = new CompOp("hi");
                    hs = new CompOp("hs");
                    equ = new CompOp("equ");
                    neu = new CompOp("neu");
                    ltu = new CompOp("ltu");
                    leu = new CompOp("leu");
                    gtu = new CompOp("gtu");
                    geu = new CompOp("geu");
                    num = new CompOp("num");
                    nan = new CompOp("nan");
                }
            }

            public class BoolOp
            {
                private string op;

                public BoolOp(string o)
                {
                    op = o;
                }

                public override string ToString()
                {
                    return op;
                }
            }

            public class BoolOps
            {
                #region static fields
                public static readonly BoolOp and;
                public static readonly BoolOp or;
                public static readonly BoolOp xor;
                #endregion

                static BoolOps()
                {
                    and = new BoolOp("and");
                    or = new BoolOp("or");
                    xor = new BoolOp("xor");
                }
            }

            public class SpaceType
            {
                private string state;

                public SpaceType(string s)
                {
                    state = s;
                }

                public override string ToString()
                {
                    return state;
                }
            }

            public class SpaceTypes
            {
                //.const, .global, .local, .param, .shared 
                #region static fields
                public static readonly SpaceType cnst;
                public static readonly SpaceType glob;
                public static readonly SpaceType loc;
                public static readonly SpaceType param;
                public static readonly SpaceType shard;
                #endregion

                static SpaceTypes()
                {
                    cnst = new SpaceType("const");
                    glob = new SpaceType("global");
                    loc = new SpaceType("local");
                    param = new SpaceType("param");
                    shard = new SpaceType("shared");
                }
            }

            public struct OpType
            {
                public OperandType opt;
                public String name;
            }

            public class OpCode
            {
                protected OperandType _opt;
                protected String _name;
                //protected byte m_s1;
                //protected OpCodeType _opc;
                //internal byte m_s2;

                //FlowControl _fc;
                //OperandType _opt;
                //internal OpCode(string name, FlowControl flcontr, OpCodeType oc)

                //internal OpCode(string name, OpCodeType oc, byte s1)
                //{
                //    this._name = name;
                //    //this._fc = flcontr;
                //    this._opc = oc;
                //    this.m_s1 = s1;
                //    //this.m_s2 = s2;
                //}
                //internal OpCode(string name, byte s1)
                internal OpCode(string name)
                {
                    this._name = name;
                    //this.m_s1 = s1;
                    this._opt = null;
                    //this._fc = flcontr;
                    //this._opc = oc;
                    //this.m_s2 = s2;
                }
                //internal OpCode(string name, byte s1, OperandType opt)
                internal OpCode(string name, OperandType opt)
                {
                    this._name = name;
                    this._opt = opt;
                    //this.m_s1 = s1;
                    //this._fc = flcontr;
                    //this._opc = oc;
                    //this.m_s2 = s2;
                }

                //
                // Summary:
                //     The name of the NVIDIA PTX instruction.
                //
                // Returns:
                //     Read-only. The name of the NVIDIA PTX instruction.
                public string Name { get { return _name; } }

                public OperandType OpType
                {
                    set { _opt = value; }
                    get { return _opt; }
                }

                public override string ToString()
                {
                    return _name;
                }
            }

            public class ALOption
            {
                private string op;

                public ALOption(string o)
                {
                    op = o;
                }

                public override string ToString()
                {
                    return op;
                }
            }
            
            public class ALOptions
            {
                #region static fields
                /// <summary>
                /// Integer operations
                /// </summary>
                public static readonly ALOption sat = new ALOption("sat"); //applied only to add/sub (32 bit) and add/sub/mul/fma/mad (64 bit)
                public static readonly ALOption cc = new ALOption("cc"); //applied only to addc/subc with 32 bit operands
                public static readonly ALOption hi = new ALOption("hi"); //applied only to mul/mad/mul24
                public static readonly ALOption lo = new ALOption("lo"); //applied only to mul/mad/mul24
                public static readonly ALOption wide = new ALOption("wide"); //applied only to mul/mad
                /// <summary>
                /// Floating point operations
                /// </summary>
                public static readonly ALOption rn = new ALOption("rn"); //applied only to add/sub/mul/fma/mad
                public static readonly ALOption rz = new ALOption("rz"); //applied only to add/sub/mul/fma/mad
                public static readonly ALOption rm = new ALOption("rm"); //applied only to add/sub/mul/fma/mad
                public static readonly ALOption rp = new ALOption("rp"); //applied only to add/sub/mul/fma/mad
                public static readonly ALOption ftz = new ALOption("ftz"); //applied only to all 
                #endregion
            }

            public class AritmLogicOpCode : OpCode
            {
                private List<ALOption> opts;

                public AritmLogicOpCode(string name, OperandType oper)
                    : base(name, oper)
                {
                    opts = new List<ALOption>(5);
                }

                public void AddOptional(ALOption opt)
                {
                    opts.Add(opt);
                }

                public override string ToString()
                {
                    StringBuilder str = new StringBuilder();
                    str.AppendFormat("{0}", _name);
                    for (int i = 0; i < opts.Count; i++)
                        str.AppendFormat(".{0}", opts[i]);
                    str.AppendFormat(".{0}", _opt.Name);
                    return str.ToString();
                }
            }

            public class CompOpCode : OpCode
            {
                private CompOp cop;
                private BoolOp bop;

                public CompOpCode(string name, OperandType oper)
                    : base(name, oper)
                {
                    cop = null;
                    bop = null;
                }

                public CompOpCode(string name, OperandType oper, CompOp c, BoolOp b)
                    : base(name, oper)
                {
                    cop = c;
                    bop = b;
                }

                public CompOp CmOp
                {
                    set { cop = value; }
                }

                public BoolOp BoOp
                {
                    set { bop = value; }
                }

                public override string ToString()
                {
                    StringBuilder str = new StringBuilder();
                    str.AppendFormat("{0}", _name);
                    if (cop != null)
                        str.AppendFormat(".{0}", cop);
                    if (bop != null)
                        str.AppendFormat(".{0}", bop);
                    str.AppendFormat(".{0}", _opt.Name);

                    return str.ToString();
                }
            }

            public class CompSetOpCode : OpCode
            {
                private CompOp cop;
                private BoolOp bop;

                public CompSetOpCode(string name, OperandType oper)
                    : base(name, oper)
                {
                    cop = null;
                    bop = null;
                }

                public CompSetOpCode(string name, OperandType oper, CompOp c, BoolOp b)
                    : base(name, oper)
                {
                    cop = c;
                    bop = b;
                }

                public CompOp CmOp
                {
                    set { cop = value; }
                }

                public BoolOp BoOp
                {
                    set { bop = value; }
                }

                public override string ToString()
                {
                    StringBuilder str = new StringBuilder();
                    str.AppendFormat("{0}", _name);
                    str.AppendFormat(".{0}", cop);
                    if (bop != null)
                        str.AppendFormat(".{0}", bop);

                    // FIX ME: PTX 1.4: set.CmpOp[.BoolOp].dtype.stype where dtype = {u,s,f}, stype = {b,u,s,f}
                    str.AppendFormat(".{0}", _opt.Name);
                    str.AppendFormat(".{0}", _opt.Name);

                    return str.ToString();
                }
            }

            public class MovOpCode : OpCode
            {
                public MovOpCode(string name, OperandType dot)
                    : base(name, dot)
                { }

                public override string ToString()
                {
                    StringBuilder str = new StringBuilder();
                    str.AppendFormat("{0}.{1}", _name, _opt.Name);
                    return str.ToString();
                }
            }

            public class LDSTOpCode : OpCode
            {
                private SpaceType st;
                private bool volatil;
                private string vect;

                public LDSTOpCode(string name, OperandType dot)
                    : base(name, dot)
                {
                    st = null;
                    volatil = false;
                    vect = null;
                }

                public LDSTOpCode(string name, OperandType dot, SpaceType s, bool vol, string v)
                    : base(name, dot)
                {
                    st = s;
                    volatil = vol;
                    vect = v;
                }

                public bool Volatile
                {
                    set { volatil = value; }
                }

                public string Vector
                {
                    set { vect = value; }
                }

                public override string ToString()
                {
                    StringBuilder str = new StringBuilder();
                    str.AppendFormat("{0}", _name);

                    if (volatil)
                        str.AppendFormat(".volatile");
                    str.AppendFormat(".{0}", st.ToString());

                    if (vect != null)
                        str.AppendFormat(".{0}", vect);

                    str.AppendFormat(".{0}", _opt.Name);
                    return str.ToString();
                }
            }

            public class ConvOpCode : OpCode
            {
                private List<string> opts;
                private OperandType sot;

                public ConvOpCode(string name, OperandType dot, OperandType s)
                    : base(name, dot)
                {
                    opts = new List<string>(2);
                    sot = s;
                }

                public OperandType DestOpType
                {
                    set { _opt = value; }
                    get { return _opt; }
                }

                public OperandType SrcOpType
                {
                    set { sot = value; }
                    get { return sot; }
                }

                public void AddOptional(string opt)
                {
                    opts.Add(opt);
                }

                public override string ToString()
                {
                    StringBuilder str = new StringBuilder();
                    str.AppendFormat("{0}", _name);
                    for (int i = 0; i < opts.Count; i++)
                        str.AppendFormat(".{0}", opts[i]);
                    str.AppendFormat(".{0}", _opt.Name);
                    str.AppendFormat(".{0}", sot.Name);
                    return str.ToString();
                }
            }

            public class FlowCtrOpCode : OpCode
            {
                private bool diverg;
                private Target target;

                public FlowCtrOpCode(string name, Target trg)
                    : base(name, null)
                {
                    diverg = true;
                    target = trg;
                }

                public override string ToString()
                {
                    StringBuilder str = new StringBuilder();
                    str.AppendFormat("{0}", _name);
                    if (!diverg)
                        str.AppendFormat(".uni");
                    if (target != null)
                        str.AppendFormat(" {0};", target.ToString());
                    return str.ToString();
                }
            }

            public class OpCodes
            {
                #region fields
                public static readonly string add;
                public static readonly string sub;
                //public static readonly OpCode addc;
                //public static readonly OpCode subc;
                public static readonly string mul;
                public static readonly string mad;
                //public static readonly OpCode mul24;
                //public static readonly OpCode mad24;
                public static readonly string sad;
                public static readonly string div;
                public static readonly string rem;
                public static readonly string abs;
                public static readonly string neg;
                public static readonly string min;
                public static readonly string max;
                public static readonly string sett;
                public static readonly string setp;
                public static readonly string selp;
                public static readonly string slct;
                public static readonly string and;
                public static readonly string or;
                public static readonly string xor;
                public static readonly string not;
                public static readonly string cnot;
                public static readonly string shl;
                public static readonly string shr;
                public static readonly string mov;
                public static readonly string ld;
                //public static readonly string ld_param;
                //public static readonly string ld_shared;
                //public static readonly string ld_global;
                public static readonly string st;
                public static readonly string cvt;
                //public OpCode tex;
                public static readonly string bra;
                public static readonly string call;
                public static readonly string ret;
                public static readonly string exit;

                public static readonly string any;
                //public OpCode bar;
                //public OpCode atom;
                //public OpCode red;
                //public OpCode vote;
                //public OpCode rcp;
                //public OpCode sqrt;
                //public OpCode rsqrt;
                //public OpCode sin;
                //public OpCode cos;
                //public OpCode lg2;
                //public OpCode ex2;
                //public OpCode trap;
                //public OpCode brkpt;

                #endregion

                static OpCodes()
                {
                    add = "add";
                    sub = "sub";
                    //public static readonly OpCode addc;
                    //public static readonly OpCode subc;
                    mul = "mul";
                    mad = "mad";
                    //public static readonly OpCode mul24;
                    //public static readonly OpCode mad24;
                    sad = "sad";
                    div = "div";
                    rem = "rem";
                    abs = "abs";
                    neg = "neg";
                    min = "min";
                    max = "max";
                    sett = "set";
                    setp = "setp";
                    selp = "selp";
                    slct = "slct";
                    and = "and";
                    or = "or";
                    xor = "xor";
                    not = "not";
                    cnot = "cnot";
                    shl = "shl";
                    shr = "shr";
                    mov = "mov";
                    ld = "ld";
                    //ld_param = "ld.param";
                    //ld_shared = "ld.shared";
                    //ld_global = "ld.global";
                    st = "st";
                    cvt = "cvt";
                    //public OpCode tex;
                    bra = "bra";
                    call = "call";
                    ret = "ret";
                    exit = "exit";

                    any = "any";
                    #region old version
                    /*
                    add = new OpCode("add", FlowControl.next, OpCodeType.aritm);
                    sub = new OpCode("sub", FlowControl.next, OpCodeType.aritm);
                    addc = new OpCode("addc", FlowControl.next, OpCodeType.aritm);
                    subc = new OpCode("subc", FlowControl.next, OpCodeType.aritm);
                    mul = new OpCode("mul", FlowControl.next, OpCodeType.aritm);
                    mad = new OpCode("mad", FlowControl.next, OpCodeType.aritm);
                    mul24 = new OpCode("mul24", FlowControl.next, OpCodeType.aritm);
                    mad24 = new OpCode("mad24", FlowControl.next, OpCodeType.aritm);
                    sad = new OpCode("sad", FlowControl.next, OpCodeType.aritm);
                    div = new OpCode("div", FlowControl.next, OpCodeType.aritm);
                    rem = new OpCode("rem", FlowControl.next, OpCodeType.aritm);
                    abs = new OpCode("abs", FlowControl.next, OpCodeType.aritm);
                    neg = new OpCode("neg", FlowControl.next, OpCodeType.aritm);
                    min = new OpCode("min", FlowControl.next, OpCodeType.aritm);
                    max = new OpCode("max", FlowControl.next, OpCodeType.aritm);
                    set = new OpCode("set", FlowControl.next, OpCodeType.compare);
                    setp = new OpCode("setp", FlowControl.next, OpCodeType.compare);
                    selp = new OpCode("selp", FlowControl.next, OpCodeType.compare);
                    slct = new OpCode("slct", FlowControl.next, OpCodeType.compare);
                    and = new OpCode("and", FlowControl.next, OpCodeType.logic);
                    or = new OpCode("or", FlowControl.next, OpCodeType.logic);
                    xor = new OpCode("xor", FlowControl.next, OpCodeType.logic);
                    not = new OpCode("not", FlowControl.next, OpCodeType.logic);
                    cnot = new OpCode("cnot", FlowControl.next, OpCodeType.logic);
                    shl = new OpCode("shl", FlowControl.next, OpCodeType.logic);
                    shr = new OpCode("shr", FlowControl.next, OpCodeType.logic);
                    mov = new OpCode("mov", FlowControl.next, OpCodeType.datamov);
                    ld = new OpCode("ld", FlowControl.next, OpCodeType.datamov);
                    st = new OpCode("st", FlowControl.next, OpCodeType.datamov);
                    cvt = new OpCode("cvt", FlowControl.next, OpCodeType.datamov);
                    tex = new OpCode("tex", FlowControl.next, OpCodeType.tex);
                    bra = new OpCode("bra", FlowControl.branch, OpCodeType.cflow);
                    call = new OpCode("call", FlowControl.call, OpCodeType.cflow);
                    ret = new OpCode("ret", FlowControl.ret, OpCodeType.cflow);
                    exit = new OpCode("exit", FlowControl.exit, OpCodeType.cflow);
                    bar = new OpCode("bar", FlowControl.next, OpCodeType.parallel);
                    atom = new OpCode("atom", FlowControl.next, OpCodeType.parallel);
                    red = new OpCode("red", FlowControl.next, OpCodeType.parallel);
                    vote = new OpCode("vote", FlowControl.next, OpCodeType.parallel);
                    rcp = new OpCode("rcp", FlowControl.next, OpCodeType.fpoint);
                    sqrt = new OpCode("sqrt", FlowControl.next, OpCodeType.fpoint);
                    rsqrt = new OpCode("rsqrt", FlowControl.next, OpCodeType.fpoint);
                    sin = new OpCode("sin", FlowControl.next, OpCodeType.fpoint);
                    cos = new OpCode("cos", FlowControl.next, OpCodeType.fpoint);
                    lg2 = new OpCode("lg2", FlowControl.next, OpCodeType.fpoint);
                    ex2 = new OpCode("ex2", FlowControl.next, OpCodeType.fpoint);
                    trap = new OpCode("trap", FlowControl.call, OpCodeType.misc);
                    brkpt = new OpCode("brkpt", FlowControl.ret, OpCodeType.misc);
                    */
                    #endregion
                }
            }

            public class Label
            {
                int id;
                public Label(int lbId) { id = lbId; }

                public int ID
                {
                    get { return id; }
                }

                public override string ToString()
                {
                    return String.Format("$Lt_0_{0}:", id);
                }
            }

            public class Target
            {
                private int id;
                public Target(int lbId) { id = lbId; }
                public override string ToString()
                {
                    return String.Format("$Lt_0_{0}", id);
                }
            }

            #endregion

            public class SharedMem
            {
                OperandType _type;
                int _align;
                string _name;
                int _count;

                public SharedMem(OperandType opt, int align, string name, int qta)
                {
                    _type = opt;
                    _align = align;
                    _name = name;
                    _count = qta;
                }

                public override string ToString()
                {
                    string res = null;
                    if (_align == 0)
                        res = String.Format(String.Format(".shared .{0} {1}[{2}];", _type.Name, _name, _count));
                    else
                        res = String.Format(String.Format(".shared .align {0} .{1} {2}[{3}];", _align, _type.Name, _name, _count));
                    return res;
                }
            }

            public class StreamArgs
            {
                /// <summary>
                /// Stream elements type
                /// </summary>
                OperandType _type;

                /// <summary>
                /// Stream name, used in the source code
                /// </summary>
                string _name;

                /// <summary>
                /// Register keys: used for stream management mapping
                /// </summary>
                int[] _regKeys;

                /// <summary>
                /// True when ld.param is already emitted into PTX code, then ld.global should be emitted
                /// </summary>
                bool _emitted;

                int _lenght;

                public StreamArgs(OperandType opt, string name)
                {
                    _type = opt;
                    _name = name;
                    _regKeys = new int[5];
                    for (int i = 0; i < 5; i++)
                        _regKeys[i] = -1;

                    _emitted = false;
                }

                public StreamArgs(OperandType opt, string name, int regReaderKey, int regIncremKey, int regIndexerKey, int regLoopBoundKey, int regLoopIndexKey, int lenght)
                {
                    _type = opt;
                    _name = name;

                    _regKeys = new int[5];
                    _regKeys[0] = regReaderKey;
                    _regKeys[1] = regIncremKey;
                    _regKeys[2] = regIndexerKey;
                    _regKeys[3] = regLoopBoundKey;
                    _regKeys[4] = regLoopIndexKey;

                    _emitted = false;

                    _lenght = lenght;
                }

                public bool Emitted
                {
                    get { return _emitted; }
                    set { _emitted = value; }
                }

                public int Lenght { get { return _lenght; } }

                public int rReaderKey
                {
                    get { return _regKeys[0]; }
                }

                public int rIncremKey
                {
                    get { return _regKeys[1]; }
                }

                public int rIndexerKey
                {
                    get { return _regKeys[2]; }
                }

                public int rLoopBoundKey  
                {
                    get { return _regKeys[3]; }
                }

                public int rLoopIndexKey
                {
                    get { return _regKeys[4]; }
                }

                public string Name
                {
                    get { return _name; }
                }

                public override string ToString()
                {
                    return String.Format(String.Format(".param .{0} {1}", _type.Name, _name)); ;
                }

                public new OperandType GetType()
                {
                    return _type;
                }
            }

            public class LocalMem
            {
                OperandType _type;
                string _name;
                int _count;

                public LocalMem(OperandType opt, string name, int num)
                {
                    _type = opt;
                    _name = name;
                    _count = num;
                }

                public override string ToString()
                {
                    string res = null;
                    if (_count == 0)
                        res = String.Format(String.Format(".reg .{0} {1};", _type.Name, _name));
                    else
                        res = String.Format(String.Format(".reg .{0} %{1}<{2}>;", _type.Name, _name, _count));
                    return res;
                }
            }

            public class Instruction
            {
                private OpCode _op;
                private Predicate _pred;
                private string[] _operands;
                private Label _label;

                public Instruction(Predicate p, Label lb, OpCode o, string[] opers)
                {
                    _op = o;
                    _pred = p;
                    _label = lb;
                    _operands = opers;
                }

                public override string ToString()
                {
                    StringBuilder instrBuilder = new StringBuilder();
                    if (_label != null)
                        instrBuilder.Append(_label.ToString());

                    instrBuilder.AppendFormat("\t");

                    if (_pred != null)
                        instrBuilder.AppendFormat("{0} ", _pred.CondToString());

                    if (_op != null)
                        instrBuilder.Append(_op.ToString());
                    //else
                    //    throw new Exception("OpCode cannot be null!");

                    if (_operands != null)
                    {
                        for (int i = 0; i < _operands.Length - 1; i++)
                            instrBuilder.AppendFormat(" {0},", _operands[i].ToString());

                        instrBuilder.AppendFormat(" {0}", _operands[_operands.Length - 1].ToString());
                        instrBuilder.Append(";");
                    }

                    return instrBuilder.ToString();
                }
            }

            public class Kernel
            {
                private string _name;
                private Dictionary<int,StreamArgs> _streams;
                private Registers _registers;
                private List<SharedMem> _sharedmem;
                private List<Instruction> _streamCode;

                private int _idRegs;
                private int _idPredicate;
                public Queue<int> idPredicate;

                public Kernel(string n)
                {
                    _name = n;
                    _streamCode = new List<Instruction>();
                    _streams = new Dictionary<int,StreamArgs>();
                    _registers = new Registers();
                    _sharedmem = new List<SharedMem>();

                    _idRegs = 1;
                    _idPredicate = 0;
                    idPredicate = new Queue<int>();
                }

                private void InternalEmit(Instruction ptx)
                {
                    _streamCode.Add(ptx);
                }

                public int AddRegID(OperandType opt)
                {
                    return _registers.Add(opt);
                }

                public bool ContainsReg(OperandType opt, int id)
                {
                    return _registers.Contains(opt, id);
                }

                public string RegString(OperandType opt, int regIndex)
                {
                    return _registers.ToString(opt, regIndex);
                }

                public int AddPredicate()
                {
                    int tmp = _idPredicate;
                    idPredicate.Enqueue(_idPredicate++);
                    return tmp;
                }

                public void AddStream(int argNr, OperandType opt, string name)
                {
                    AddStream(argNr, opt, name);
                }

                public void AddStream(int argNr, OperandType opt, string name, int regReaderKey, int regIncremKey, int regIndexerKey, int regLoopBoundKey, int regLoopIndexKey, int lenght)
                {
                    if (_streams.ContainsKey(argNr))
                        throw new Exception(String.Format("Cannot insert the stream {0} since there is another stream with the same key already inserted",argNr));
                    
                    _streams.Add(argNr, new StreamArgs(opt, name, regReaderKey, regIncremKey, regIndexerKey, regLoopBoundKey, regLoopIndexKey, lenght));
                }

                public StreamArgs GetStream(int index)
                {
                    return _streams[index];
                }

                public int GetPredicate()
                {
                    if (idPredicate.Count > 0)
                        return idPredicate.Dequeue();
                    else
                        return -1;
                }

                public string Name
                {
                    get { return _name; }
                }

                public virtual void AddSharedMem(OperandType opt, string name, int qta)
                {
                    _sharedmem.Add(new SharedMem(opt, 4, name, qta));
                }

                public virtual void AddSharedMem(OperandType opt, int align, string name, int qta)
                {
                    _sharedmem.Add(new SharedMem(opt, align, name, qta));
                }

                public virtual void EmitLabel(Label lab)
                {
                    _streamCode.Add(new Instruction(null, lab, null, null));
                }

                public virtual void EmitAR(string opcode, OperandType oper, string[] args, List<ALOption> opts)
                {
                    EmitAR(null, null, opcode, oper, args, opts);
                }

                public virtual void EmitAR(Label lab, string opcode, OperandType oper, string[] args, List<ALOption> opts)
                {
                    EmitAR(lab, null, opcode, oper, args, opts);
                }

                public virtual void EmitAR(Predicate pred, string opcode, OperandType oper, string[] args, List<ALOption> opts)
                {
                    EmitAR(null, pred, opcode, oper, args, opts);
                }

                public virtual void EmitAR(Label lab, Predicate pred, string opcode, OperandType oper, string[] args, List<ALOption> opts)
                {
                    //this.EnsureCapacity(3);
                    AritmLogicOpCode opc = new AritmLogicOpCode(opcode, oper);
                    if (opts != null)
                    {
                        for (int i = 0; i < opts.Count; i++)
                            opc.AddOptional(opts[i]);
                    }
                    Instruction instr = new Instruction(pred, lab, opc, args);
                    this.InternalEmit(instr);
                }

                public virtual void EmitCMP(string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    EmitCMP(null, null, opcode, oper, co, bo, args);
                }

                public virtual void EmitCMP(Label lab, string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    EmitCMP(lab, null, opcode, oper, co, bo, args);
                }

                public virtual void EmitCMP(Predicate pred, string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    EmitCMP(null, pred, opcode, oper, co, bo, args);
                }

                public virtual void EmitCMP(Label lab, Predicate pred, string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    //this.EnsureCapacity(3);
                    CompOpCode opc = new CompOpCode(opcode, oper, co, bo);
                    Instruction instr = new Instruction(pred, lab, opc, args);
                    this.InternalEmit(instr);
                }

                public virtual void EmitSetCMP(string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    EmitSetCMP(null, null, opcode, oper, co, bo, args);
                }

                public virtual void EmitSetCMP(Label lab, string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    EmitSetCMP(lab, null, opcode, oper, co, bo, args);
                }

                public virtual void EmitSetCMP(Predicate pred, string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    EmitSetCMP(null, pred, opcode, oper, co, bo, args);
                }

                public virtual void EmitSetCMP(Label lab, Predicate pred, string opcode, OperandType oper, CompOp co, BoolOp bo, string[] args)
                {
                    //this.EnsureCapacity(3);
                    CompSetOpCode opc = new CompSetOpCode(opcode, oper, co, bo);
                    Instruction instr = new Instruction(pred, lab, opc, args);
                    this.InternalEmit(instr);
                }

                public virtual void EmitMV(string opcode, OperandType dest, string[] args)
                {
                    EmitMV(null, null, opcode, dest, args);
                }

                public virtual void EmitMV(Label lab, string opcode, OperandType dest, string[] args)
                {
                    EmitMV(lab, null, opcode, dest, args);
                }

                public virtual void EmitMV(Predicate pred, string opcode, OperandType dest, string[] args)
                {
                    EmitMV(null, pred, opcode, dest, args);
                }

                public virtual void EmitMV(Label lab, Predicate pred, string opcode, OperandType dest, string[] args)
                {
                    //this.EnsureCapacity(3);
                    MovOpCode opc = new MovOpCode(opcode, dest);
                    Instruction instr = new Instruction(pred, lab, opc, args);
                    this.InternalEmit(instr);
                }

                public virtual void EmitLS(string opcode, OperandType dest, SpaceType st, bool volat, string vect, string[] args)
                {
                    EmitLS(null, null, opcode, dest, st, volat, vect, args);
                }

                public virtual void EmitLS(Label lab, string opcode, OperandType dest, SpaceType st, bool volat, string vect, string[] args)
                {
                    EmitLS(lab, null, opcode, dest, st, volat, vect, args);
                }

                public virtual void EmitLS(Predicate pred, string opcode, OperandType dest, SpaceType st, bool volat, string vect, string[] args)
                {
                    EmitLS(null, pred, opcode, dest, st, volat, vect, args);
                }

                public virtual void EmitLS(Label lab, Predicate pred, string opcode, OperandType dest, SpaceType st, bool volat, string vect, string[] args)
                {
                    //this.EnsureCapacity(3);
                    LDSTOpCode opc = new LDSTOpCode(opcode, dest, st, volat, vect);
                    Instruction instr = new Instruction(pred, lab, opc, args);
                    this.InternalEmit(instr);
                }

                public virtual void EmitCV(string opcode, OperandType dest, OperandType src, string[] args)
                {
                    EmitCV(null, null, opcode, dest, src, args);
                }

                public virtual void EmitCV(Label lab, string opcode, OperandType dest, OperandType src, string[] args)
                {
                    EmitCV(lab, null, opcode, dest, src, args);
                }

                public virtual void EmitCV(Predicate pred, string opcode, OperandType dest, OperandType src, string[] args)
                {
                    EmitCV(null, pred, opcode, dest, src, args);
                }

                public virtual void EmitCV(Label lab, Predicate pred, string opcode, OperandType dest, OperandType src, string[] args)
                {
                    //this.EnsureCapacity(3);
                    ConvOpCode opc = new ConvOpCode(opcode, dest, src);
                    Instruction instr = new Instruction(pred, lab, opc, args);
                    this.InternalEmit(instr);
                }

                public virtual void EmitFC(string opcode, Target target)
                {
                    EmitFC(null, null, opcode, target);
                }

                public virtual void EmitFC(Label lab, string opcode, Target target)
                {
                    EmitFC(lab, null, opcode, target);
                }

                public virtual void EmitFC(Predicate pred, string opcode, Target target)
                {
                    EmitFC(null, pred, opcode, target);
                }

                public virtual void EmitFC(Label lab, Predicate pred, string opcode, Target target)
                {
                    //this.EnsureCapacity(3);
                    FlowCtrOpCode opc = new FlowCtrOpCode(opcode, target);
                    Instruction instr = new Instruction(pred, null, opc, null);
                    this.InternalEmit(instr);
                }

                public override string ToString()
                {
                    StringBuilder strbld = new StringBuilder();
                    strbld.AppendFormat("\t.entry {0} (", _name);
                    strbld.AppendLine();
                    for (int index = 0; index < _streams.Count - 1; index++)
                    {
                        strbld.AppendFormat("\t\t{0},", _streams[index].ToString());
                        strbld.AppendLine();
                    }
                    strbld.AppendFormat("\t\t{0} )", _streams[_streams.Count - 1].ToString());
                    strbld.AppendLine();
                    strbld.AppendLine("\t{");
                    
                    // Registers
                    Dictionary<OperandType, List<int>>.Enumerator iter = _registers.GetEnumerator();
                    while (iter.MoveNext())
                        strbld.AppendLine(String.Format("\t.reg .{0} {1}<{2}>;", iter.Current.Key.Name, _registers.opTypeName(iter.Current.Key), _registers.opTypeCount(iter.Current.Key) + 1));

                    // Predicates
                    strbld.AppendLine(String.Format("\t.reg .pred %p<{0}>;", _idPredicate + 1));
                    // Shared memory
                    for (int index = 0; index < _sharedmem.Count; index++)
                        strbld.AppendLine(_sharedmem[index].ToString());
                    // PTX code
                    foreach (Instruction code in _streamCode)
                        strbld.AppendLine(code.ToString());

                    strbld.AppendLine("\texit;");
                    strbld.AppendLine("\t}");

                    return strbld.ToString();
                }
            }

            static public class TypesMapping
            {
                static public OperandType Convert(System.Type tp, params SystemType[] st)
                {
                    SystemType sysT = SystemType.x86;

                    if (tp == null)
                        return null;

                    if (st.Length > 0)
                        sysT = st[0];

                    if (tp.Equals(typeof(System.Boolean)))
                        return OperandTypes.b8;

                    if (tp.Equals(typeof(System.Int16)))
                        if(sysT != SystemType.x64)
                            return OperandTypes.s16;
                        else
                            return OperandTypes.s64;

                    if (tp.Equals(typeof(System.Int32)))
                        if (sysT != SystemType.x64)
                            return OperandTypes.s32;
                        else
                            return OperandTypes.s64;

                    if (tp.Equals(typeof(System.Int64)))
                        return OperandTypes.s64;

                    if (tp.Equals(typeof(System.UInt16)))
                        if (sysT != SystemType.x64)
                            return OperandTypes.u16;
                        else
                            return OperandTypes.u64;

                    if (tp.Equals(typeof(System.UInt32)))
                        if (sysT != SystemType.x64)
                            return OperandTypes.u32;
                        else
                            return OperandTypes.u64;

                    if (tp.Equals(typeof(System.UInt64)))
                        return OperandTypes.u64;

                    if (tp.Equals(typeof(System.Single)))
                        return OperandTypes.f32;

                    if (tp.Equals(typeof(System.Double)))
                        return OperandTypes.f64;

                    if (tp.Equals(typeof(System.Byte)))
                        return OperandTypes.b8;        

                    if (tp.Equals(typeof(OutputStream<>)))
                        return OperandTypes.param;

                    if (tp.Equals(typeof(InputStream<>)))
                        return OperandTypes.param;

                    if (tp.Equals(typeof(Target)))
                        return OperandTypes.target;

                    if (tp.Equals(typeof(MethodDesc)))
                        return OperandTypes.method;

                    if (tp.Equals(typeof(System.Collections.IEnumerator)))
                        return OperandTypes.u64;

                    if (tp.Equals(typeof(void)))
                        return OperandTypes.method;

                    //if (tp.Equals(typeof(System.String)))
                    //    return  OperandTypes.;
                    throw new ArgumentException();
                }
            }

            #region Compiler building blocks: Node and OperandInfo

            public class OperandInfo : IComparable
            {
                public OperandType optp;
                public System.Reflection.Emit.OpCode opc;
                public OpCode ptxOpc;
                public Object Value;

                public OperandInfo(System.Reflection.Emit.OpCode oc, OperandType ot, Object val)
                {
                    opc = oc;
                    optp = ot;
                    ptxOpc = MapOpcodeCILPTX(oc, ot);
                    Value = val;
                }

                public OperandInfo(string oc, OperandType ot, Object val)
                {
                    opc = System.Reflection.Emit.OpCodes.Nop;
                    optp = ot;
                    ptxOpc = new OpCode(oc, ot); //MapOpcodeCILPTX(oc, ot);
                    Value = val;
                }

                private OpCode MapOpcodeCILPTX(System.Reflection.Emit.OpCode ocp, OperandType ot)
                {
                    OPCODES op = (OPCODES)((ocp.Value < 0 ? 256 : 0) + (ocp.Value & 0xFF));
                    switch (op)
                    {
                        case OPCODES.Add:
                            return new OpCode(OpCodes.add, ot);
                        case OPCODES.Sub:
                            return new OpCode(OpCodes.sub, ot);
                        case OPCODES.Div:
                            return new OpCode(OpCodes.div, ot);
                        case OPCODES.And:
                            return new OpCode(OpCodes.and, ot);
                        case OPCODES.Mul:
                            return new OpCode(OpCodes.mul, ot);
                        case OPCODES.Br:
                        case OPCODES.Br_S:
                        case OPCODES.Brtrue:
                        case OPCODES.Brfalse:
                        case OPCODES.Brfalse_S:
                        case OPCODES.Brtrue_S:
                            return new OpCode(OpCodes.bra, ot);
                        case OPCODES.Call:
                        case OPCODES.Callvirt:
                            return new OpCode(OpCodes.call, ot);
                        case OPCODES.Ldarg:
                            return new OpCode(OpCodes.ld, ot);
                        case OPCODES.Starg:
                            return new OpCode(OpCodes.st, ot);
                        case OPCODES.Ldloc:
                            return new OpCode(OpCodes.mov, ot);
                        case OPCODES.Ldfld:
                            return new OpCode(OpCodes.ld, ot);
                        case OPCODES.Stloc:
                            return new OpCode(OpCodes.st, ot);
                        case OPCODES.Cgt:
                        case OPCODES.Ceq:
                        case OPCODES.Clt:
                            return new OpCode(OpCodes.sett, ot);
                        case OPCODES.Beq:
                        case OPCODES.Bge:
                        case OPCODES.Bgt:
                        case OPCODES.Ble:
                        case OPCODES.Blt:
                            return new OpCode(OpCodes.bra, ot);
                        case OPCODES.Ldc_I4:
                        case OPCODES.Ldc_I8:
                        case OPCODES.Ldc_R4:
                        case OPCODES.Ldc_R8:
                            return new OpCode(OpCodes.mov, ot);
                        case OPCODES.Box:
                        case OPCODES.Unbox:
                        case OPCODES.Unbox_Any:
                            return new OpCode(OpCodes.any, ot);
                        case OPCODES.Nop:
                            return new OpCode(OpCodes.any, ot);
                        default:
                            throw new NotImplementedException();
                    }
                }

                public OpCode PTXOPC
                {
                    get { return ptxOpc; }
                }

                public override bool Equals(object obj)
                {
                    return ((obj == null) && (Value == null));
                }

                public bool IsNull()
                {
                    return (Value == null);
                }

                public OPCODES Kind()
                {
                    return (OPCODES)((opc.Value < 0 ? 256 : 0) + (opc.Value & 0xFF));
                }

                public static bool operator !=(OperandInfo a, OperandInfo b)
                {
                    return !(a.Equals(b));
                }

                public static bool operator ==(OperandInfo a, OperandInfo b)
                {
                    return a.Equals(b);
                }

                public override int GetHashCode()
                {
                    return this.opc.GetHashCode();
                }

                public int CompareTo(object value)
                {
                    if (value == null)
                    {
                        return 1;
                    }
                    if (!(value is OperandInfo))
                    {
                        throw new ArgumentException();
                    }
                    return Compare(this, (OperandInfo)value);
                }

                public virtual int Compare(OperandInfo a, OperandInfo b)
                {
                    if (a == null)
                        return -1;

                    if (b == null)
                        return 1;

                    if (a == b)
                        return 0;

                    throw new NotSupportedException();
                }
            }

            /// <summary>
            /// Base class for tree nodes
            /// </summary>
            public class Node
            {
                protected int id;
                protected OperandInfo opcode;
                protected int regId;
                protected Label label;
                protected Kernel kers;
                protected List<Node> nodes;

                public Node(Kernel krs, Label lab, int identifier, OperandInfo opc)
                {
                    kers = krs;
                    id = identifier;
                    opcode = opc;
                    regId = -1;
                    label = lab;
                    nodes = new List<Node>();
                }

                public int GetID()
                {
                    return id;
                }

                public int GetRegID()
                {
                    return regId;
                }

                public OperandType GetOPType()
                {
                    return opcode.optp;
                }

                public OperandInfo Info
                {
                    get { return opcode; }
                }

                public void SetType(OperandType tp)
                {
                    opcode.optp = tp;
                }

                public Label GetLabel()
                {
                    return label;
                }

                public void AddNode(Node nd)
                {
                    nodes.Add(nd);
                }

                public virtual string Scan()
                {
                    string reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);
                    
                    reg = kers.RegString(opcode.optp, ind);
                    
                    if (nodes.Count != 0)
                    {
                        for(int i = 0; i < nodes.Count; ++i)
                            nodes[i].Scan();
                    }

                    if (label != null)
                        kers.EmitLabel(label);
                    
                    ///TODO: put here a FP64 converter
                    return reg;
                }
            }

            /// <summary>
            /// NOP  instructions
            /// </summary>
            public class NodeNOP : Node
            {
                public NodeNOP(Kernel krs, Label lab, int identifier, OperandInfo opc)
                    : base(krs, lab, identifier, opc) { }

                public override string Scan()
                {
                    if (label != null)
                        kers.EmitLabel(label);
                    return "";
                }
            }

            public class NodeARG : Node
            {
                string name;
                public NodeARG(Kernel krs, Label lab, int identifier, OperandInfo opc, string nm)
                    : base(krs, lab, identifier, opc)
                {
                    name = nm;
                }

                public override string Scan()
                {
                    if (label != null)
                        kers.EmitLabel(label);

                    return id.ToString(); 
                }
            }

            /// <summary>
            /// Node for numerical constant
            /// </summary>
            public class NodeConst : Node // stack behaviour: pop = 0, push = 1 
            {
                public NodeConst( Kernel krs, Label lab, OperandInfo opc)
                    : base( krs, lab, -1, opc)
                {
                }

                public override string Scan()
                {
                    if (label != null)
                        kers.EmitLabel(label);
                    return String.Format("{0}", opcode.Value);
                }
            }

            /// <summary>
            /// Ldloc / Stloc
            /// </summary>
            public class NodeLSloc : Node
            {
                public NodeLSloc(Kernel krs, Label lab, int identifier, OperandInfo opcd)
                    : base(krs, lab, identifier, opcd) { }

                public override string Scan()
                {
                    string[] results = new string[2];

                    if( nodes.Count > 0)
                    {
                        results[1] = nodes[0].Scan();

                        if ((nodes[0].Info.ptxOpc.Name == OpCodes.setp) ||
                             (nodes[0].Info.ptxOpc.Name == OpCodes.any))  
                            return results[1];
                    }

                    string reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);

                    if (nodes[0].Info.ptxOpc.Name == OpCodes.ld)  
                            return results[1];

                    results[0] = reg;

                    int pred = kers.GetPredicate();
                    if (pred >= 0)
                        kers.EmitMV(label, new Predicate(true, pred), OpCodes.mov, opcode.optp, results);
                    else
                        kers.EmitMV(label, OpCodes.mov, opcode.optp, results);

                    return results[0];
                }
            }

            /// <summary>
            /// Arithemical - logical instructions
            /// </summary>
            public class NodeCALC : Node
            {
                List<ALOption> opts;

                public NodeCALC(Kernel krs, Label lab, OperandInfo opcd)
                    : base(krs, lab, -1, opcd)
                {
                    opts = null;
                }

                public NodeCALC(Kernel krs, Label lab, int identifier, OperandInfo opcd)
                    : base( krs, lab, identifier, opcd)
                {
                    opts = null;
                }

                public NodeCALC(Kernel krs, Label lab, int identifier, OperandInfo opcd, List<ALOption> options)
                    : base(krs, lab, identifier, opcd)
                {
                    opts = options;
                }

                public override string Scan()
                {
                    string[] results = new string[nodes.Count + 1];
                    OperandType optype = nodes[0].GetOPType() > nodes[1].GetOPType() ? nodes[0].GetOPType() : nodes[1].GetOPType();

                    // TODO: insert a conversion instruction whenever required

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);
                    
                    results[0] = reg;

                    if (label != null)
                        kers.EmitLabel(label);

                    int pred = kers.GetPredicate();
                    if (pred >= 0)
                        kers.EmitAR(new Predicate(true, pred), opcode.ptxOpc.Name, optype, results, opts);
                    else
                        kers.EmitAR(opcode.ptxOpc.Name, optype, results, opts);
                    return results[0];
                }
            }

            /// <summary>
            /// Compare operations
            /// </summary>
            public class NodeCMP : Node
            {
                protected CompOp comp;
                protected BoolOp boolop;

                public NodeCMP(Kernel krs, Label lab, OperandInfo opcd)
                    : base( krs, lab, -1, opcd)
                {
                    comp = null;
                    boolop = null;
                }

                public NodeCMP(Kernel krs, Label lab, OperandInfo opcd, CompOp compare)
                    : base( krs, lab, -1, opcd)
                {
                    comp = compare;
                    boolop = null;
                }

                public NodeCMP(Kernel krs, Label lab, OperandInfo opcd, CompOp cop, BoolOp bop)
                    : base( krs, lab, -1, opcd)
                {
                    comp = cop;
                    boolop = bop;
                }

                public CompOp CompareOp
                {
                    set { comp = value; }
                }

                public BoolOp BooleanOp
                {
                    set { boolop = value; }
                }

                public override string Scan()
                {
                    //  setp.le.s32 %p1, %r2, %r3;   
                    string[] results = new string[nodes.Count + 1];

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    regId = kers.AddPredicate();
                    results[0] = String.Format("%p{0}", regId);

                    OperandType optype = nodes[0].GetOPType() > nodes[1].GetOPType() ? nodes[0].GetOPType() : nodes[1].GetOPType();

                    kers.EmitCMP(label, opcode.ptxOpc.Name, optype, comp, boolop, results);
                    return results[0];
                }
            }

            /// <summary>
            /// Compare operations
            /// </summary>
            public class NodeSetCMP : NodeCMP
            {
                public NodeSetCMP( Kernel krs, Label lab, OperandInfo opcd)
                    : base(krs, lab, opcd) { }

                public NodeSetCMP(Kernel krs, Label lab, OperandInfo opcd, CompOp compare)
                    : base(krs, lab, opcd, compare) { }

                public NodeSetCMP(Kernel krs, Label lab, OperandInfo opcd, CompOp cop, BoolOp bop)
                    : base(krs, lab, opcd, cop, bop) { }

                public override string Scan()
                {
                    //  setp.le.s32 %p1, %r2, %r3;   
                    string[] results = new string[nodes.Count + 1];

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    OperandType optype = nodes[0].GetOPType() > nodes[1].GetOPType() ? nodes[0].GetOPType() : nodes[1].GetOPType();

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);
                    results[0] = reg;

                    kers.EmitSetCMP(label, opcode.ptxOpc.Name, optype, comp, boolop, results);
                    return results[0];
                }
            }

            /// <summary>
            /// Branch instructions
            /// </summary>
            public class NodeBRA : Node
            {
                protected Target target;
                protected Predicate predicate;

                public NodeBRA(Kernel krs, Label lab, OperandInfo opcd)
                    : base( krs, lab, -1, opcd)
                {
                    target = null;
                    predicate = null;
                }

                public NodeBRA(Kernel krs, Label lab, OperandInfo opcd, Target trg)
                    : base(krs, lab, -1, opcd)
                {
                    target = trg;
                    predicate = null;
                }

                public NodeBRA(Kernel krs, Label lab, OperandInfo opcd, Target trg, Predicate pred)
                    : base( krs, lab, -1, opcd)
                {
                    target = trg;
                    predicate = pred;
                }

                public Target Target
                {
                    set { target = value; }
                    get { return target; }
                }

                public Predicate Predic
                {
                    set { predicate = value; }
                    get { return predicate; }
                }

                public override string Scan()
                {
                    //  @%p1 bra $Lt_0_1282;
                    string[] results = new string[nodes.Count];

                    for (int n = 0; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    int pred = kers.GetPredicate();
                    if ((pred >= 0) && (predicate == null))
                        kers.EmitFC(label, new Predicate(true, pred), opcode.ptxOpc.Name, target);
                    else
                        kers.EmitFC(label, predicate, opcode.ptxOpc.Name, target);

                    return null;
                }
            }

            /// <summary>
            /// Flow control instructions
            /// </summary>
            public class NodeFCBRA : NodeCMP
            {
                protected Target target;
                protected Predicate predicate;

                public NodeFCBRA( Kernel krs, Label lab, OperandInfo opcd, CompOp co)
                    : base( krs, lab, opcd, co)
                {
                    target = null;
                    predicate = null;
                    //regId = kers.AddRegID();
                }

                public NodeFCBRA( Kernel krs, Label lab, OperandInfo opcd, CompOp co, BoolOp bo)
                    : base( krs, lab, opcd, co, bo)
                {
                    target = null;
                    predicate = null;
                    //regId = kers.AddRegID();
                }

                public NodeFCBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, CompOp co)
                    : base( krs, lab, opcd, co)
                {
                    target = trg;
                    predicate = null;
                    //regId = kers.AddRegID();
                }

                public NodeFCBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, CompOp co, BoolOp bo)
                    : base( krs, lab, opcd, co, bo)
                {
                    target = trg;
                    predicate = null;
                    //regId = kers.AddRegID();
                }

                public NodeFCBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, Predicate pred, CompOp co)
                    : base( krs, lab, opcd, co)
                {
                    target = trg;
                    predicate = pred;
                    //regId = kers.AddRegID();
                }

                public NodeFCBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, Predicate pred, CompOp co, BoolOp bo)
                    : base( krs, lab, opcd, co, bo)
                {
                    target = trg;
                    predicate = pred;
                    //regId = kers.AddRegID();
                }

                public Target Target
                {
                    set { target = value; }
                    get { return target; }
                }

                public Predicate Predic
                {
                    set { predicate = value; }
                    get { return predicate; }
                }

                public override string Scan()
                {
                    //  @%p1 bra $Lt_0_1282;
                    string[] results = new string[nodes.Count + 1];

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    results[0] = predicate.Name;

                    OperandType optype = nodes[0].GetOPType() > nodes[1].GetOPType() ? nodes[0].GetOPType() : nodes[1].GetOPType();

                    kers.EmitCMP(OpCodes.setp, optype, comp, boolop, results);
                    kers.EmitFC(label, predicate, opcode.ptxOpc.Name, target);

                    return null;
                }
            }

            /// <summary>
            /// Branch instructions
            /// </summary>
            public class NodeCONDBRA : NodeFCBRA
            {
                string compValue;
                public NodeCONDBRA(Kernel krs, Label lab, OperandInfo opcd, CompOp co, string cmpValue)
                    : base( krs, lab, opcd, co) { compValue = cmpValue; }

                public NodeCONDBRA( Kernel krs, Label lab, OperandInfo opcd, CompOp co, BoolOp bo, string cmpValue)
                    : base( krs, lab, opcd, co, bo) { compValue = cmpValue; }

                public NodeCONDBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, CompOp co, string cmpValue)
                    : base( krs, lab, opcd, trg, co) { compValue = cmpValue; }

                public NodeCONDBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, CompOp co, BoolOp bo, string cmpValue)
                    : base( krs, lab, opcd, trg, co, bo) { compValue = cmpValue; }

                public NodeCONDBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, Predicate pred, CompOp co, string cmpValue)
                    : base( krs, lab, opcd, trg, pred, co) { compValue = cmpValue; }

                public NodeCONDBRA( Kernel krs, Label lab, OperandInfo opcd, Target trg, Predicate pred, CompOp co, BoolOp bo, string cmpValue)
                    : base( krs, lab, opcd, trg, pred, co, bo) { compValue = cmpValue; }

                public override string Scan()
                {
                    string[] results1 = new string[3];
                    string[] results = new string[nodes.Count + 1];
                    OperandType optype = opcode.optp;

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);
                    results[0] = reg;

                    results1[2] = compValue;
                    results1[1] = results[0];
                    results1[0] = predicate.Name;

                    kers.EmitCMP(OpCodes.setp, optype, comp, boolop, results1);
                    kers.EmitFC(label, predicate, opcode.ptxOpc.Name, target);

                    return null;
                }
            }

            public class NodeSTglob : Node
            {
                public NodeSTglob( Kernel krs, Label lab, int identifier, OperandInfo opcd)
                    : base( krs, lab, identifier, opcd) { }

                public override string Scan()
                {
                    // st.global into output param
                    string[] results = new string[nodes.Count];
                    
                    OperandType optype = nodes[0].GetOPType() > nodes[1].GetOPType() ? nodes[0].GetOPType() : nodes[1].GetOPType();

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);

                    results[0] = String.Format("[{0}+0]", reg);

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[n].Scan();

                    int pred = kers.GetPredicate();
                    if (pred >= 0)
                        kers.EmitLS(label, new Predicate(true, pred), opcode.ptxOpc.Name, optype, SpaceTypes.glob, false, null, results);
                    else
                        kers.EmitLS(label, opcode.ptxOpc.Name, optype, SpaceTypes.glob, false, null, results);

                    return results[0];
                }
            }

            public class NodeLDglob : Node
            {
                SpaceType _space;

                public NodeLDglob( Kernel krs, SpaceType st, Label lab, int identifier, OperandInfo opcd)
                    : base( krs, lab, identifier, opcd) 
                {
                    _space = st;
                }

                public override string Scan()
                {
                    string[] results = new string[2];

                    results[0] = nodes[0].Scan();
                    OperandType optype = nodes[0].GetOPType();

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);
                    results[1] = reg;

                    if (_space == SpaceTypes.param)
                        results[1] = String.Format("[{0}]", results[1]);
                    else
                        results[1] = String.Format("[{0}+0]", results[1]);

                    int pred = kers.GetPredicate();
                    if (pred >= 0)
                        kers.EmitLS(label, new Predicate(true, pred), opcode.ptxOpc.Name, optype, _space, false, null, results);
                    else
                        kers.EmitLS(label, opcode.ptxOpc.Name, optype, _space, false, null, results);

                    return results[0];
                }
            }

            public class NodeLSshared : Node
            {
                public NodeLSshared( Kernel krs, Label lab, int identifier, OperandInfo opcd)
                    : base( krs, lab, identifier, opcd) { }

                public override string Scan()
                {
                    string[] results = new string[nodes.Count + 1];
                    OperandType optype = nodes[0].GetOPType() > nodes[1].GetOPType() ? nodes[0].GetOPType() : nodes[1].GetOPType();

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);

                    results[0] = reg;

                    int pred = kers.GetPredicate();
                    if (pred >= 0)
                        kers.EmitLS(label, new Predicate(true, pred), opcode.ptxOpc.Name, optype, SpaceTypes.shard, false, null, results);
                    else
                        kers.EmitLS(label, opcode.ptxOpc.Name, optype, SpaceTypes.shard, false, null, results);

                    return results[0];
                }
            }

            public class NodeCALL : Node
            {
                SpaceType spacet;

                public NodeCALL( Kernel krs, Label lab, int identifier, OperandInfo opcd,
                                 SpaceType spt)
                    : base(krs, lab, identifier, opcd)
                {
                    spacet = spt;
                }

                public override string Scan()
                {
                    string[] results = new string[nodes.Count + 1];
                    OperandType optype = null;
                    if (nodes.Count > 1)
                        optype = nodes[0].GetOPType() > nodes[1].GetOPType() ? nodes[0].GetOPType() : nodes[1].GetOPType();
                    else
                        optype = nodes[0].GetOPType();

                    for (int n = 1; n < results.Length; n++)
                        results[n] = nodes[results.Length - n - 1].Scan();

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);

                    results[0] = reg;

                    kers.EmitLS(label, opcode.ptxOpc.Name, optype, spacet, false, null, results);
                    return results[0];
                }
            }

            public class NodeCVT : Node
            {
                OperandInfo _dest;

                public NodeCVT( Kernel krs, Label lab, OperandInfo src, OperandInfo dest)
                    : base( krs, lab, -1, src)
                {
                    _dest = dest;
                }

                public override string Scan()
                {
                    string[] results;

                    if (nodes.Count == 0)
                        results = new string[1];
                    else
                    {
                        results = new string[2];
                        results[1] = nodes[0].Scan();
                    }

                    String reg = null;
                    int ind = id;

                    if (!kers.ContainsReg(opcode.optp, id))
                        ind = kers.AddRegID(opcode.optp);

                    reg = kers.RegString(opcode.optp, ind);

                    results[0] = reg;

                    int pred = kers.GetPredicate();
                    if (pred >= 0)
                        kers.EmitCV(label, new Predicate(true, pred), OpCodes.cvt, _dest.optp, opcode.optp ,results);
                    else
                        kers.EmitCV(label, OpCodes.cvt, _dest.optp, opcode.optp, results);

                    return results[0];
                }
            }

            #endregion

            public class KernelFile
            {
                string code;
                string fileName;
                string gen_registers;
                string version;
                string target;

                public KernelFile(string kernel_path)
                {
                    version = "\t.version 1.4";
                    target = "\t.target sm_12, map_f64_to_f32";

                    //StringBuilder stbRegs = new StringBuilder();
                    //stbRegs.AppendLine("\t.reg .u32 %ra<17>;");
                    //stbRegs.AppendLine("\t.reg .u64 %rda<17>;");
                    //stbRegs.AppendLine("\t.reg .f32 %fa<17>;");
                    //stbRegs.AppendLine("\t.reg .f64 %fda<17>;");
                    //stbRegs.AppendLine("\t.reg .u32 %rv<5>;");
                    //stbRegs.AppendLine("\t.reg .u64 %rdv<5>;");
                    //stbRegs.AppendLine("\t.reg .f32 %fv<5>;");
                    //stbRegs.AppendLine("\t.reg .f64 %fdv<5>;");
                    gen_registers = "";//stbRegs.ToString();

                    fileName = kernel_path;
                    code = null;
                }

                public KernelFile(string kernel_path, string capability)
                {
                    //capabil = capability;
                    version = "\t.version 1.4";
                    target = String.Format("\t.target sm_{0}, map_f64_to_f32", capability);

                    //StringBuilder stbRegs = new StringBuilder();
                    //stbRegs.AppendLine("\t.reg .u32 %ra<17>;");
                    //stbRegs.AppendLine("\t.reg .u64 %rda<17>;");
                    //stbRegs.AppendLine("\t.reg .f32 %fa<17>;");
                    //stbRegs.AppendLine("\t.reg .f64 %fda<17>;");
                    //stbRegs.AppendLine("\t.reg .u32 %rv<5>;");
                    //stbRegs.AppendLine("\t.reg .u64 %rdv<5>;");
                    //stbRegs.AppendLine("\t.reg .f32 %fv<5>;");
                    //stbRegs.AppendLine("\t.reg .f64 %fdv<5>;");
                    gen_registers = "";//stbRegs.ToString();

                    fileName = kernel_path;
                    code = null;
                }

                public string Version
                {
                    get { return version; }
                    set { version = value; }
                }

                public string Target
                {
                    get { return target; }
                    set { target = value; }
                }

                public string Body
                {
                    set { code = value; }
                }

                public string FileName
                {
                    get { return fileName; }
                }

                public void Write()
                {
                    if (code == null)
                        throw new ArgumentNullException();

                    StringBuilder stb = new StringBuilder();
                    stb.AppendFormat(version);
                    stb.AppendLine();
                    stb.AppendFormat(target);
                    stb.AppendLine();
                    stb.AppendLine("");
                    stb.AppendLine(gen_registers);

                    StreamWriter sw = File.CreateText(fileName);
                    sw.WriteLine(stb.ToString());
                    sw.WriteLine(code);
                    sw.Flush();
                    sw.Close();
                }
            }

            public class Registers
            {
                private const string _prefixName = "%r";
                private Dictionary<OperandType, List<int>> _regs;

                public Registers()
                {
                    _regs = new Dictionary<OperandType, List<int>>();
                }
                
                public int Add(OperandType opt)
                {
                    if (!_regs.ContainsKey(opt))
                        _regs.Add(opt,new List<int>());

                    int index = _regs[opt].Count + 1;
                    _regs[opt].Add(index);
                    
                    return index;
                }

                public int opTypeCount(OperandType t)
                {
                    return _regs[t].Count;
                }

                public bool Contains(OperandType t, int id)
                {
                    return _regs[t].Contains(id);
                }

                public string opTypeName(OperandType t)
                {
                    return String.Format("{0}{1}", _prefixName, t.Id4RegisterNaming);
                }

                public string ToString(OperandType t, int id)
                {
                    return String.Format("{0}{1}{2}", _prefixName, t.Id4RegisterNaming, _regs[t][id-1]);
                }

                public Dictionary<OperandType, List<int>>.Enumerator GetEnumerator()
                {
                    return _regs.GetEnumerator();
                }
            }

            /// <summary>
            /// TODO: check GetCurret and SetCurrent methods, wrong ptx generation. It seems branch generation correct, but not the same for 
            ///       iterator management. Hints: check increment by NodeCALC and MoveNext 
            /// </summary>
            public class MetaCUDAKernel : MetaKernel
            {
                // Generated code
                private Kernel _ptxCode;

                // Increment of stream indexer
                private const int increm = 4;

                public MetaCUDAKernel(Object worker, MethodInfo minf, SystemType st)
                    : base(worker, minf, st)
                {
                    _ptxCode = new Kernel(_meth.Name);
                }

                public MetaCUDAKernel(Object worker, MethodInfo minf, int ncores, SystemType st)
                    : base(worker, minf, ncores, st)
                {
                    _ptxCode = new Kernel(_meth.Name);
                }

                /// <summary>
                /// ILCodeScan
                /// </summary>
                /// <param name="op"></param>
                /// <param name="cil"></param>
                /// <param name="NodesStack"></param>
                /// <param name="targets"></param>
                /// <param name="labe"></param>
                /// <param name="predicateCounter"></param>
                protected void ILCodeScan(OPCODES op, ILInstruction cil, int[] stackBehav, Stack<Node> NodesStack,
                                        Dictionary<CLIFile.Target, Target> targets, Label labe, ref int predicateCounter)
                {
                    switch (op)
                    {
                        case OPCODES.Stloc:
                            {
                                Node tmp = NodesStack.Pop();
                                OperandType loctype = TypesMapping.Convert(_localMemory[cil.par.iv].LocalType, sType);
                                OperandType optype = tmp.GetOPType() > loctype ? tmp.GetOPType() : loctype;
                                tmp.SetType(optype);

                                NodeLSloc sto = new NodeLSloc(_ptxCode, labe, tmp.GetID(), new OperandInfo(cil.op, optype, tmp.GetID()));
                                sto.AddNode(tmp);
                                NodesStack.Push(sto);
                            }
                            break;
                        case OPCODES.Ldarg:
                            {
                                if (cil.par.iv > 0)
                                {
                                    int index = cil.par.iv - 1;
                                    // Current stream info
                                    StreamArgs sa = _ptxCode.GetStream(index);
                                    NodesStack.Push(new NodeARG(_ptxCode, labe, sa.rLoopIndexKey, new OperandInfo(cil.op,
                                                                                          TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                                          index), _inputs[index].Name));
                                }
                            }
                            break;
                        case OPCODES.Ldloc:
                        case OPCODES.Ldloc_S:
                            NodesStack.Push(new NodeConst(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                                      TypesMapping.Convert(_localMemory[cil.par.iv].LocalType, sType),
                                                                                      String.Format("loc{0}", cil.par.iv))));
                            break;
                        case OPCODES.Ldc_I4:
                            NodesStack.Push(new NodeConst(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                          TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                          cil.par.iv)));
                            break;
                        case OPCODES.Ldc_I8:
                            NodesStack.Push(new NodeConst(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                          TypesMapping.Convert(cil.par.lv.GetType(), sType),
                                                                          cil.par.lv)));
                            break;
                        case OPCODES.Ldc_R4:
                            NodesStack.Push(new NodeConst(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                          TypesMapping.Convert(cil.par.fv.GetType(), sType),
                                                                          cil.par.fv)));
                            break;
                        case OPCODES.Ldc_R8:
                            NodesStack.Push(new NodeConst(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                          TypesMapping.Convert(cil.par.dv.GetType(), sType),
                                                                          cil.par.dv)));
                            break;
                        case OPCODES.Conv_I:
                        case OPCODES.Conv_I1:
                        case OPCODES.Conv_I2:
                        case OPCODES.Conv_I4:
                            {
                                Node tmp = NodesStack.Pop();
                                NodeCVT nbr = new NodeCVT(  _ptxCode, labe, tmp.Info, new OperandInfo(cil.op,
                                                                           TypesMapping.Convert(typeof(int), sType), null));

                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                            }
                            break;

                        case OPCODES.Conv_I8:
                            {
                                Node tmp = NodesStack.Pop();
                                NodeCVT nbr = new NodeCVT(  _ptxCode, labe, tmp.Info, new OperandInfo(cil.op,
                                                                           TypesMapping.Convert(typeof(long), sType), null));

                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                            }
                            break;
                        case OPCODES.Conv_R4:
                            {
                                Node tmp = NodesStack.Pop();
                                NodeCVT nbr = new NodeCVT(  _ptxCode, labe, tmp.Info, new OperandInfo(cil.op,
                                                                           TypesMapping.Convert(typeof(float), sType), null));

                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                            }
                            break;
                        case OPCODES.Conv_R8:
                            {
                                Node tmp = NodesStack.Pop();
                                NodeCVT nbr = new NodeCVT(  _ptxCode, labe, tmp.Info, new OperandInfo(cil.op,
                                                                           TypesMapping.Convert(typeof(double), sType), null));

                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                            }
                            break;
                        case OPCODES.Call:
                        case OPCODES.Calli:
                        case OPCODES.Callvirt:
                            {
                                /// TODO: which type is the actual parameter?
                                MethodDesc trg = (MethodDesc)_cursor.Instr.ResolveParameter(_clifile);
                                int numArgs = trg.Signature.Count + (trg.Signature.HasThis ? 1 : 0); // + 1 is the instance var

                                Node[] tmp = new Node[stackBehav[1]];
                                for (int i = 0; i < stackBehav[1]; i++)
                                    tmp[stackBehav[1] - i - 1] = NodesStack.Pop();

                                if (tmp[0].GetType() == typeof(NodeARG))
                                {
                                    int idx = (int)tmp[0].Info.Value;
                                    Type ptm1 = _inputs[idx].ParameterType;
                                    Type ptm2 = typeof(StreamDefs.InputStream<>);
                                    bool cond = false;
                                    OperandType opt = tmp[0].GetOPType();

                                    // Current stream info
                                    StreamArgs sa = _ptxCode.GetStream(idx);

                                    if (ptm1.IsGenericType)
                                        cond = (ptm1.Name.CompareTo("InputStream`1") == 0);
                                    else
                                        cond = ptm1 == ptm2;
                                    
                                    if ( cond )
                                    {
                                        if ((trg.Name.CompareTo("get_Item") == 0) || (trg.Name.CompareTo("get_Current") == 0))
                                        {
                                            SpaceType st = SpaceTypes.glob;

                                            NodeLDglob nldg = new NodeLDglob(  _ptxCode, st, labe, sa.rIndexerKey, new OperandInfo(System.Reflection.Emit.OpCodes.Ldarg,
                                                                tmp[0].GetOPType(), null));

                                            // Both return an element pushed onto the stack
                                            int regId = _ptxCode.AddRegID(sa.GetType());

                                            nldg.AddNode(new Node(_ptxCode, labe, regId, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, sa.GetType(), null)));
                                            
                                            NodesStack.Push(nldg);
                                        }
                                        else if (trg.Name.CompareTo("MoveNext") == 0) 
                                        {
                                            // add.s32 	%r_loop_index, %r_loop_index, val_inc;
                                            // mov.s32 	%r_loop_bound, streamLength;
                                            // setp.ge.s32 	%px, %r_loop_index, %r_loop_bound;
                                            NodeCALC nbr = new NodeCALC(  _ptxCode, labe, sa.rLoopIndexKey, new OperandInfo(System.Reflection.Emit.OpCodes.Add,
                                                                                                    opt, null));

                                            nbr.AddNode(new NodeConst(  _ptxCode, labe, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, 1)));
                                            nbr.AddNode(new Node(  _ptxCode, tmp[0].GetLabel(), sa.rLoopIndexKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt,null)));
                                            
                                            NodeLSloc nsc = new NodeLSloc(  _ptxCode, labe, sa.rLoopBoundKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null));
                                            nsc.AddNode(new Node(  _ptxCode, labe, sa.rLoopBoundKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null)));

                                            NodeCMP ncmp = new NodeCMP(  _ptxCode, labe, new OperandInfo(OpCodes.setp, opt, null), CompOps.ne);
                                            ncmp.AddNode(nsc);
                                            ncmp.AddNode(nbr);
                                            
                                            NodesStack.Push(ncmp);
                                        }
                                        else
                                        {
                                            NodeCALL nbr = new NodeCALL(  _ptxCode, labe, tmp[0].GetID(), new OperandInfo(System.Reflection.Emit.OpCodes.Ldarg,
                                                                opt, trg), SpaceTypes.param);

                                            for (int i = 0; i < stackBehav[1]; i++)
                                                nbr.AddNode(tmp[i]);

                                            NodesStack.Push(nbr);
                                        }
                                    }
                                    //else if (_inputs[idx].ParameterType == typeof(StreamDefs.OutputStream<>))
                                    if (ptm1.Name.CompareTo("OutputStream`1") == 0)
                                    {
                                        if ((trg.Name.CompareTo("set_Item") == 0) || (trg.Name.CompareTo("set_Current") == 0))
                                        {
                                            NodeSTglob nbr = new NodeSTglob(  _ptxCode, labe, sa.rReaderKey, new OperandInfo(System.Reflection.Emit.OpCodes.Starg,
                                                               opt, null));

                                            for (int i = 0; i < stackBehav[1]; i++)
                                                nbr.AddNode(tmp[i]);

                                            NodesStack.Push(nbr);
                                        }
                                        else if (trg.Name.CompareTo("MoveNext") == 0)
                                        {
                                            Label currlabe = tmp[0].GetLabel();
                                            // Increment indexer over output stream
                                            NodeConst ncst1 = new NodeConst(  _ptxCode, null, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, increm));
                                            Node ncst2 = new Node(  _ptxCode, null, sa.rReaderKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null));

                                            NodeCALC nadd = new NodeCALC(  _ptxCode, currlabe, sa.rReaderKey, new OperandInfo(System.Reflection.Emit.OpCodes.Add,
                                                                            opt, sa.rReaderKey));
                                            nadd.AddNode(ncst1);
                                            nadd.AddNode(ncst2);
                                            
                                            NodeCALC nbr = new NodeCALC(  _ptxCode, null, sa.rLoopIndexKey, new OperandInfo(System.Reflection.Emit.OpCodes.Add,
                                                                                                    opt, null));

                                            Node temp = new Node(  _ptxCode, null, sa.rLoopIndexKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null));
                                            temp.AddNode(nadd);

                                            nbr.AddNode(new NodeConst(  _ptxCode, null, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, 1)));
                                            nbr.AddNode(temp);

                                            NodeLSloc nsc = new NodeLSloc(  _ptxCode, null, sa.rLoopBoundKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null));
                                            nsc.AddNode(new Node(  _ptxCode, labe, sa.rLoopBoundKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null)));

                                            NodeCMP ncmp = new NodeCMP(  _ptxCode, null, new OperandInfo(OpCodes.setp, opt, null), CompOps.ne);
                                            ncmp.AddNode(nsc);
                                            ncmp.AddNode(nbr);

                                            NodesStack.Push(ncmp);
                                        }
                                        else
                                            throw new MethodAccessException();
                                    }
                                    else if (_inputs[idx].ParameterType == typeof(StreamDefs.Stream))
                                        throw new NotImplementedException();

                                }
                                else //static methods
                                {
                                    /// TODO: 2. gestire la call threadId()
                                    if (trg.Name.CompareTo("threadId") == 0)
                                    {
                                        //cvt.s32.u16 	%r9, %tid.x;
                                        NodeCVT nbr = new NodeCVT(  _ptxCode, labe,
                                                                    new OperandInfo(System.Reflection.Emit.OpCodes.Nop,
                                                                        OperandTypes.u16, null),
                                                                    new OperandInfo(System.Reflection.Emit.OpCodes.Conv_I,
                                                                        OperandTypes.s32, null));

                                        NodesStack.Push(nbr);
                                    }
                                }
                            }
                            break;
                        case OPCODES.Br:
                            {
                                // Branch target
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                Target to = null;
                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                    {
                                        to = new Target((int)trg.Position);
                                        break;
                                    }

                                NodesStack.Push(new NodeBRA(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                              TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                              trg), to));
                            }
                            break;
                        case OPCODES.Brtrue_S:
                        case OPCODES.Brtrue:
                            {
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                Target to = null;

                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                    {
                                        to = new Target((int)trg.Position);
                                        break;
                                    }

                                NodeBRA nbr = new NodeBRA(  _ptxCode, labe, new OperandInfo(cil.op, TypesMapping.Convert(cil.par.iv.GetType(), sType), null), to, new Predicate(true, predicateCounter++));

                                Node node = NodesStack.Pop();
                                nbr.AddNode(node);
                                NodesStack.Push(nbr);
                                //bran.Scan();
                            }
                            break;
                        case OPCODES.Brfalse_S:
                        case OPCODES.Brfalse:
                            {
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                Target to = null;

                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                    {
                                        to = new Target((int)trg.Position);
                                        break;
                                    }

                                NodeCONDBRA nbr = new NodeCONDBRA(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                  trg), to, new Predicate(true, predicateCounter++),
                                                                  CompOps.eq, "0x00000000");

                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //bran.Scan();
                            }
                            break;
                        case OPCODES.Beq:
                            {
                                Target to = null;
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                    {
                                        to = new Target((int)trg.Position);
                                        break;
                                    }

                                NodeFCBRA nbr = new NodeFCBRA(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                  trg), to, new Predicate(true, predicateCounter++),
                                                                  CompOps.eq);
                                nbr.AddNode(NodesStack.Pop());
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //nbr.Scan();
                            }
                            break;
                        case OPCODES.Bge:
                            {
                                Target to = null;
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                        to = new Target((int)trg.Position);

                                NodeFCBRA nbr = new NodeFCBRA(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                  trg), to, new Predicate(true, predicateCounter++),
                                                                  CompOps.ge);
                                nbr.AddNode(NodesStack.Pop());
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //nbr.Scan();
                            }
                            break;
                        case OPCODES.Bgt:
                            {
                                Target to = null;
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                        to = new Target((int)trg.Position);

                                NodeFCBRA nbr = new NodeFCBRA(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                  trg), to, new Predicate(true, predicateCounter++),
                                                                  CompOps.gt);
                                nbr.AddNode(NodesStack.Pop());
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //nbr.Scan();
                            }
                            break;
                        case OPCODES.Ble:
                            {
                                Target to = null;
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                        to = new Target((int)trg.Position);

                                NodeFCBRA nbr = new NodeFCBRA(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                  trg), to, new Predicate(true, predicateCounter++),
                                                                  CompOps.le);
                                nbr.AddNode(NodesStack.Pop());
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                ///nbr.Scan();
                            }
                            break;
                        case OPCODES.Blt:
                            {
                                Target to = null;
                                CLIFile.Target trg = (CLIFile.Target)_cursor.Instr.ResolveParameter(_clifile);
                                for (int y = 0; y < _cursor.Targets.Length; y++)
                                    if (_cursor.Targets[y].CompareTo(trg) == 0)
                                        to = new Target((int)trg.Position);

                                NodeFCBRA nbr = new NodeFCBRA(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  TypesMapping.Convert(cil.par.iv.GetType(), sType),
                                                                  trg), to, new Predicate(true, predicateCounter++),
                                                                  CompOps.lt);
                                nbr.AddNode(NodesStack.Pop());
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //nbr.Scan();
                            }
                            break;
                        case OPCODES.Ceq:
                            {
                                Node node = NodesStack.Pop();
                                NodeSetCMP nbr = new NodeSetCMP(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  node.GetOPType(),
                                                                  cil.op.Value),
                                                                  CompOps.eq);
                                nbr.AddNode(node);
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //nbr.Scan();
                            }
                            break;
                        case OPCODES.Cgt:
                            {
                                Node node = NodesStack.Pop();
                                NodeSetCMP nbr = new NodeSetCMP(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  node.GetOPType(),
                                                                  cil.op.Value),
                                                                  CompOps.gt);
                                nbr.AddNode(node);
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //nbr.Scan();
                            }
                            break;
                        case OPCODES.Clt:
                            {
                                Node node = NodesStack.Pop();
                                NodeSetCMP nbr = new NodeSetCMP(  _ptxCode, labe, new OperandInfo(cil.op,
                                                                  node.GetOPType(),
                                                                  cil.op.Value),
                                                                  CompOps.lt);
                                nbr.AddNode(node);
                                nbr.AddNode(NodesStack.Pop());
                                NodesStack.Push(nbr);
                                //nbr.Scan();
                            }
                            break;
                        case OPCODES.Mul:
                            {
                                Node node1 = NodesStack.Pop();
                                Node node2 = NodesStack.Pop();
                                List<ALOption> options = new List<ALOption>(1);
                                options.Add( ALOptions.lo );

                                OperandType optype = node1.GetOPType() > node2.GetOPType() ? node1.GetOPType() : node2.GetOPType();

                                NodeCALC root = new NodeCALC(_ptxCode, labe, (int)_cursor.Position, new OperandInfo(cil.op,
                                                                                                    optype,
                                                                                                    cil.op.Value), options);
                                root.AddNode(node1);
                                root.AddNode(node2);
                                NodesStack.Push(root);
                            }
                            break;
                        case OPCODES.Add:
                        case OPCODES.Sub:
                        case OPCODES.Div:
                            {
                                Node node1 = NodesStack.Pop();
                                Node node2 = NodesStack.Pop();

                                OperandType optype = node1.GetOPType() > node2.GetOPType() ? node1.GetOPType() : node2.GetOPType();

                                NodeCALC root = new NodeCALC(_ptxCode, labe, (int)_cursor.Position, new OperandInfo(cil.op,
                                                                                                    optype,
                                                                                                    cil.op.Value));
                                root.AddNode(node1);
                                root.AddNode(node2);
                                NodesStack.Push(root);
                            }
                            break;
                        case OPCODES.Nop:
                            NodesStack.Push(new NodeNOP(  _ptxCode, labe, -1, null));
                            break;
                        //case OPCODES.Unbox:
                        //case OPCODES.Unbox_Any:
                        //case OPCODES.Box:
                        //    {
                        //        //if(NodesStack.Count>0)
                        //        //{
                        //        //    Node tmp = NodesStack.Peek();
                        //        //    NodesStack.Push(new NodeConst(  _ptxCode, labe, tmp.GetID(), new OperandInfo(tmp.Info.opc, null, tmp.GetRegID())));
                        //        //}
                        //        //else
                        //        //    NodesStack.Push(new NodeConst(  _ptxCode, labe, null, new OperandInfo(cil.op, null, "")));
                        //    }
                        //    break;
                    }
                }

                protected void initRegister4Streams(StreamDefs.Stream args, int index, string name)
                {
                    /// Registers IDs required to map the first stream access pattern
                    int regIdReader = 0;
                    int regIdIncrem = 0;
                    int regIdIndexer = 0;

                    /// Registers IDs required to map the loop control flow over streams pattern
                    int regIdLoopBound = 0; // Loop bounded to stream length
                    int regIdLoopIndex = 0;

                    string[] paramsIndex = new string[2];
                    string[] paramsLength = new string[2];
                    string[] paramsStream = new string[2];

                    OperandType opt = TypesMapping.Convert(args.GetType(), sType);

                    regIdReader = _ptxCode.AddRegID(opt);
                    string rreaderKey = _ptxCode.RegString(opt, regIdReader);

                    regIdLoopIndex = _ptxCode.AddRegID(opt);
                    string rloopIdxKey = _ptxCode.RegString(opt, regIdLoopIndex);

                    regIdIncrem = _ptxCode.AddRegID(opt);
                    
                    regIdIndexer = _ptxCode.AddRegID(opt);

                    // Insert auxiliary registers to PTX code immediately after declarations
                    // used to iterate over streams.
                    paramsIndex[1] = "0"; // init register to 0, it will point to the first element in the stream
                    paramsIndex[0] = rloopIdxKey;
                    _ptxCode.EmitMV(OpCodes.mov, opt, paramsIndex);

                    // Insert auxiliary registers to PTX code immediately after declarations
                    // used as bound of iterations over streams.
                    // For each stream:
                    // 1. determine its length
                    // 2. set the value of a register to that length.
                    //Type ttg1 = tt.GetGenericTypeDefinition();
                    int argLength = args.Count();

                    regIdLoopBound = _ptxCode.AddRegID(opt);
                    string rloopBndKey = _ptxCode.RegString(opt, regIdLoopBound);

                    paramsLength[1] = argLength.ToString(); // init register to 0, it will point to the first element in the stream
                    paramsLength[0] = rloopBndKey;
                    _ptxCode.EmitMV(OpCodes.mov, opt, paramsLength);

                    // ld.param.u64 %r_read, [_stream_];
                    paramsStream[1] = String.Format("[{0}]", name); // init register to 0, it will point to the first element in the stream
                    paramsStream[0] = rreaderKey;
                    _ptxCode.EmitLS(null, null, OpCodes.ld, opt, SpaceTypes.param, false, null, paramsStream);
                    
                    /// inp is used as key in the Dictionary of streams
                    /// Later, during the Kernel body scanning, this key will be used to map loops over each stream
                    /// Important: when Key is 0 means the first parameter of the Kernel method.
                    _ptxCode.AddStream(index, opt, name, regIdReader, regIdIncrem, regIdIndexer, regIdLoopBound, regIdLoopIndex, args.Count());
                }

                private void initIndexes(int streamIndex)
                {
                    StreamArgs sa = _ptxCode.GetStream(streamIndex);
                    OperandType opt = sa.GetType();
                    //########################################################################################################################
                    // mul.lo.u64 %r_next, %r_loop_indexT, 4;
                    NodeConst ncst = new NodeConst(_ptxCode, null, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, increm));
                    Node nind = new Node(_ptxCode, null, sa.rLoopIndexKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null));

                    NodeCALC nmul = new NodeCALC(_ptxCode, null, sa.rIncremKey, new OperandInfo(System.Reflection.Emit.OpCodes.Mul,
                                        opt, sa.rIncremKey), new List<ALOption>(1) { ALOptions.lo });

                    nmul.AddNode(ncst);
                    nmul.AddNode(nind);

                    Node nldg = new Node(_ptxCode, null, sa.rReaderKey, new OperandInfo(System.Reflection.Emit.OpCodes.Nop, opt, null));

                    // add.u64 %r_indexer, %r_read, %r_next;
                    NodeCALC nadd = new NodeCALC(_ptxCode, null, sa.rIndexerKey, new OperandInfo(System.Reflection.Emit.OpCodes.Add,
                                        opt, sa.rIndexerKey));

                    nadd.AddNode(nmul);
                    nadd.AddNode(nldg);

                    nadd.Scan();
                    //########################################################################################################################
                }

                /// <summary>
                /// TODO: 
                /// </summary>
                public string EvalBody(StreamDefs.Stream outArg, params StreamDefs.Stream[] inputArgs)
                {
                    int globalStackHeight = 0;

                    Stack<Node> NodesStack = new Stack<Node>();
                    Dictionary<CLIFile.Target, Target> targets = new Dictionary<CLIFile.Target, Target>();
                    int predicateCounter = 0;

                    _clifile = _cursor.FileReader;
                    _cursor.TrackTargets();
                    _cursor.TrackStack(_args, _locvars, _meth.GetMethodBody().MaxStackSize);

                    #region shared memory
                    //if (_sharedMemory.Count > 0)
                    //{
                    //    Dictionary<FieldInfo, MemoryAccessType>.Enumerator iterShr = _sharedMemory.GetEnumerator();
                    //    while (iterShr.MoveNext())
                    //    {
                    //        int regId = 0;
                    //        string[] results = new string[2];
                    //        FieldInfo fino = iterShr.Current.Key;
                    //        OperandType opt = TypesMapping.Convert(fino.FieldType, sType);

                    //        regId = _ptxCode.AddRegID(opt);
                    //        string name = iterShr.Current.Key.Name;

                    //        int qta = 1;
                    //        if (fino.FieldType.IsArray)
                    //            qta = ((Array)(fino.GetValue(instance))).Length;

                    //        _ptxCode.AddSharedMem(opt, name, qta);
                    //    }
                    //}
                    #endregion

                    #region Input/Output stream

                    int streamIndex = 0;

                    /// Kernel input arguments evaluation
                    if ((inputArgs != null) && (inputArgs.Length > 0))
                    {
                        for (streamIndex = 0; streamIndex < inputArgs.Length; streamIndex++)
                        {
                            initRegister4Streams(inputArgs[streamIndex], streamIndex, _inputs[streamIndex].Name);
                            initIndexes(streamIndex);
                        }
                    }

                    /// Kernel output arguments evaluation
                    if (outArg != null)
                        initRegister4Streams(outArg, streamIndex, _inputs[streamIndex].Name);

                    #endregion

                    while (_cursor.Next())
                    {
                        Label labe = null;
                        ILInstruction cil = ILInstruction.Normalize(_cursor.Instr);
                        OPCODES op = (OPCODES)((cil.op.Value < 0 ? 256 : 0) + (cil.op.Value & 0xFF));

                        int[] stackBehav = GetBehaviour(op);

                        // Branch label
                        // se la label dell'istruzione corrente e` target di un branch 
                        // allora devo aggiungerla nella generazione dell'istruzione PTX corrente
                        if (_cursor.Label != null)
                        {
                            for (int y = 0; y < _cursor.Targets.Length; y++)
                                if (_cursor.Targets[y].CompareTo(_cursor.Label) == 0)
                                    labe = new Label((int)_cursor.Label.Position);
                        }
                        
                        ILCodeScan(op, cil, stackBehav, NodesStack,   targets, labe, ref predicateCounter);

                        globalStackHeight += stackBehav[0] - stackBehav[1];

                        Console.WriteLine("op:{0} - NSTKHGT:{1} - globStackHGT:{2}", cil.op.Name, NodesStack.Count, globalStackHeight);

                        if (globalStackHeight == 0)
                        {
                            if (NodesStack.Count > 0)
                            {
                                Node root = NodesStack.Pop();
                                string result = root.Scan();
                            }
                        }
                    }
                    Console.WriteLine(_ptxCode.ToString());
                    return _ptxCode.ToString();
                }

                public String Name
                {
                    get { return _ptxCode.Name; }
                }
            }

            public class MetaCUDACompute : MetaCompute
            {
                private KernelFile file;

                public MetaCUDACompute(Object worker, CLIFileReader clireader, SystemType st)
                    : base(worker, clireader, st)
                {
                    Type tp = worker.GetType();
                    _fields = tp.GetFields(BindingFlags.Instance | BindingFlags.Static | BindingFlags.Public);
                    
                    file = new KernelFile(String.Format("{0}_001.ptx", tp.Name), "12");
                }

                public new MetaCUDAKernel GetKernel(string methodName)
                {
                    return (MetaCUDAKernel)_kernels[methodName];
                }

                public KernelFile ptxFile
                {
                    get { return file; }
                }

                public override void AddKernel(MethodInfo kernel)
                {
                    if (kernel == null)
                        throw new ArgumentOutOfRangeException();

                    _kernels.Add(kernel.Name, new MetaCUDAKernel(instance, kernel, sType));
                }

                public void AddKernel(MethodInfo kernel, int num_cores)
                {
                    if ((kernel == null) || (num_cores < 0))
                        throw new ArgumentOutOfRangeException();

                    _kernels.Add(kernel.Name, new MetaCUDAKernel(instance, kernel, num_cores, sType));
                }
            }

            public class MetaCUDACompiler
            {
                /// <summary>
                /// Determine whether there is a GPU card installed in the current system or not.
                /// If yes, it should evaluate which GPU to use between ATI and NVIDIA.
                /// TODO: actually it statically selects NVIDIA 
                /// </summary>
                /// <param name="instanceType"></param>
                /// <param name="clifile"></param>
                /// <param name="compute"></param>
                /// <returns></returns>
                static private MetaCompute GetGPULib(Type instanceType, CLIFileReader clifile)
                {
                    return new MetaCUDACompute(instanceType, clifile, MetaCompiler.GetSystemType());
                }

                static public MetaCompute FindKernels(Type instanceType, MethodInfo compute)
                {
                    //Type typeExe = instance.GetType();
                    string tmp = instanceType.Assembly.CodeBase;
                    CLIFileReader clifile = CLIFileReader.Open(tmp);

                    object[] attribs = compute.GetCustomAttributes(false);
                    if (attribs.Length == 0)
                        return null; //continue
                    int KernelAttribIndex = Array.FindIndex<Object>(attribs,
                                                             delegate(Object o)
                                                             {
                                                                 return o.GetType().Equals(typeof(KernelsAttribute));
                                                             });
                    if (KernelAttribIndex == -1)
                        return null; //continue

                    MetaCompute computation = GetGPULib(instanceType, clifile);
                    computation.AddKernel(compute);

                    return computation;
                }

                static public MetaCompute FindKernels(Type instanceType)
                {
                    string tmp = instanceType.Assembly.Location;
                    CLIFileReader clifile = CLIFileReader.Open(tmp);
                    if (clifile == null)
                        throw new NullReferenceException();

                    MethodInfo[] methods = instanceType.GetMethods(BindingFlags.Instance | BindingFlags.Public | BindingFlags.DeclaredOnly);
                    MetaCompute computation = GetGPULib(instanceType, clifile);
                    
                    for (int me = 0; me < methods.Length; me++)
                    {
                        object[] attribs = methods[me].GetCustomAttributes(false);
                        if (attribs.Length == 0)
                            continue;

                        int KernelAttribIndex = Array.FindIndex<Object>(attribs,
                                                                 delegate(Object o)
                                                                 {
                                                                     return o.GetType().Equals(typeof(KernelsAttribute));
                                                                 });
                        if (KernelAttribIndex == -1)
                            continue;

                        computation.AddKernel(methods[me]);
                    }

                    return computation;
                }

            }
        }

        //namespace AMD
        //{
        //    //public class AMDMetaComp : MetaCompute
        //    //{

        //    //}

        //    //public class Metacomp
        //    //{
        //    //    private Type data; 

        //    //    public Metacomp(Type inputType) 
        //    //    {
        //    //        data = inputType;
        //    //    }

        //    //private virtual MethodInfo findMethod(string methodName)
        //    //{ 
        //    //    return data.GetMethod(methodName);
        //    //}
        //}
    }
}